EXSWHTEC-224 - Test cases ID clean up and documentation for Module Management

catch/hipTestMain/config/config_amd_linux

@@ -234,6 +234,21 @@
         "Unit_atomicExch_system_Positive_Host_And_Peer_GPUs - unsigned long long",
         "Unit_atomicExch_system_Positive_Host_And_Peer_GPUs - float",
         "Unit_atomicExch_system_Positive_Host_And_Peer_GPUs - double",
+        "Note: Test disabled due to defect - EXSWHTEC-151",
+        "Unit_hipModuleLoad_Negative_Load_From_A_File_That_Is_Not_A_Module",
+        "Note: Following two tests disabled due to defect - EXSWHTEC-153",
+        "Unit_hipModuleLoadData_Negative_Image_Is_An_Empty_String",
+        "Unit_hipModuleLoadDataEx_Negative_Image_Is_An_Empty_String",
+        "Note: Test disabled due to defect - EXSWHTEC-163",
+        "Unit_hipModuleGetGlobal_Negative_Hmod_Is_Nullptr",
+        "Note: Test disabled due to defect - EXSWHTEC-164",
+        "Unit_hipModuleGetGlobal_Negative_Name_Is_Empty_String",
+        "Note: Test disabled due to defect - EXSWHTEC-165",
+        "Unit_hipModuleGetGlobal_Negative_Dptr_And_Bytes_Are_Nullptr",
+        "Note: Test disabled due to defect - EXSWHTEC-166",
+        "Unit_hipModuleGetTexRef_Negative_Hmod_Is_Nullptr",
+        "Note: Test disabled due to defect - EXSWHTEC-167",
+        "Unit_hipModuleGetTexRef_Negative_Name_Is_Empty_String",
         "=== SWDEV-439004: Below tests failing randomly in CQE staging ===",
         "Unit_hipGLGetDevices_Positive_Basic",
         "Unit_hipGLGetDevices_Positive_Parameters",

catch/hipTestMain/config/config_amd_windows

@@ -291,6 +291,23 @@
         "Unit_atomicExch_system_Positive_Host_And_Peer_GPUs - unsigned long long",
         "Unit_atomicExch_system_Positive_Host_And_Peer_GPUs - float",
         "Unit_atomicExch_system_Positive_Host_And_Peer_GPUs - double",
+        "Note: Test disabled due to defect - EXSWHTEC-151",
+        "Unit_hipModuleLoad_Negative_Load_From_A_File_That_Is_Not_A_Module",
+        "Note: Test disabled due to defect - EXSWHTEC-152",
+        "Unit_hipModuleUnload_Negative_Module_Is_Nullptr",
+        "Note: Following two tests disabled due to defect - EXSWHTEC-153",
+        "Unit_hipModuleLoadData_Negative_Image_Is_An_Empty_String",
+        "Unit_hipModuleLoadDataEx_Negative_Image_Is_An_Empty_String",
+        "Note: Test disabled due to defect - EXSWHTEC-163",
+        "Unit_hipModuleGetGlobal_Negative_Hmod_Is_Nullptr",
+        "Note: Test disabled due to defect - EXSWHTEC-164",
+        "Unit_hipModuleGetGlobal_Negative_Name_Is_Empty_String",
+        "Note: Test disabled due to defect - EXSWHTEC-165",
+        "Unit_hipModuleGetGlobal_Negative_Dptr_And_Bytes_Are_Nullptr",
+        "Note: Test disabled due to defect - EXSWHTEC-166",
+        "Unit_hipModuleGetTexRef_Negative_Hmod_Is_Nullptr",
+        "Note: Test disabled due to defect - EXSWHTEC-167",
+        "Unit_hipModuleGetTexRef_Negative_Name_Is_Empty_String",
         "=== SWDEV-435667: Below tests failing randomly in stress test on 08/12/23 ===",
         "Unit_hipMemPoolSetAccess_Negative_Parameters",
         "Unit_hipMallocMipmappedArray_Negative_Parameters",

catch/include/hip_test_defgroups.hh

@@ -131,6 +131,13 @@ THE SOFTWARE.
  * @}
  */
 
+/**
+ * @defgroup ModuleTest Module Management
+ * @{
+ * This section describes tests for the module management functions of HIP runtime API.
+ * @}
+ */
+
 /**
  * @defgroup ShflTest warp shuffle function Management
  * @{

catch/unit/CMakeLists.txt

@@ -36,11 +36,11 @@ add_subdirectory(compiler)
 add_subdirectory(errorHandling)
 add_subdirectory(cooperativeGrps)
 add_subdirectory(context)
+add_subdirectory(module)
 add_subdirectory(device_memory)
 add_subdirectory(warp)
 add_subdirectory(dynamicLoading)
 add_subdirectory(g++)
-add_subdirectory(module)
 add_subdirectory(channelDescriptor)
 add_subdirectory(executionControl)
 add_subdirectory(vector_types)

catch/unit/module/get_tex_ref_module.cc

@@ -25,4 +25,4 @@ THE SOFTWARE.
 
 texture<float, 2> tex;
 
-#endif // CUDA_VERSION < CUDA_12000
+#endif // CUDA_VERSION < CUDA_12000

catch/unit/module/hipExtModuleLaunchKernel.cc

@@ -236,6 +236,41 @@ TEST_CASE("Unit_hipExtModuleLaunchKernel_UniformWorkGroup") {
   HIP_CHECK(hipModuleUnload(Module));
 }
 
+/**
+ * @addtogroup hipExtModuleLaunchKernel hipExtModuleLaunchKernel
+ * @{
+ * @ingroup ModuleTest
+ * `hipExtModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX, 
+ * uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ,
+ * uint32_t localWorkSizeX, uint32_t localWorkSizeY,
+ * uint32_t localWorkSizeZ, size_t sharedMemBytes,
+ * hipStream_t hStream, void** kernelParams, void** extra,
+ * hipEvent_t startEvent = nullptr,
+ * hipEvent_t stopEvent = nullptr, uint32_t flags = 0)` -
+ * Launches kernel with parameters and shared memory on stream with arguments passed
+ * to kernel params or extra arguments.
+ */
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Launch kernels in different basic scenarios:
+ *    -# When kernel is launched with no arguments
+ *      - Expected output: return `hipSuccess`
+ *    -# When kernel is launched with arguments using `kernelParams`
+ *      - Expected output: return `hipSuccess`
+ *    -# When kernel is launched with arguments using `extra`
+ *      - Expected output: return `hipSuccess`
+ *    -# When kernel is launched as timed and with events
+ *      - Expected output: return `hipSuccess`
+ * Test source
+ * ------------------------
+ *  - unit/module/hipExtModuleLaunchKernel.cc
+ * Test requirements
+ * ------------------------
+ *  - Platform specific (AMD)
+ *  - HIP_VERSION >= 5.2
+ */
 TEST_CASE("Unit_hipExtModuleLaunchKernel_Positive_Basic") {
   ModuleLaunchKernelPositiveBasic<hipExtModuleLaunchKernel>();
 
@@ -258,6 +293,33 @@ TEST_CASE("Unit_hipExtModuleLaunchKernel_Positive_Basic") {
   }
 }
 
+/**
+ * Test Description
+ * ------------------------
+ *  - Launches kernel with different valid arguments:
+ *    -# When gridDimX == maxGridDimX
+ *      - Expected output: return `hipSuccess`
+ *    -# When gridDimY == maxGridDimY
+ *      - Expected output: return `hipSuccess`
+ *    -# When gridDimZ == maxGridDimZ
+ *      - Expected output: return `hipSuccess`
+ *    -# When blockDimX == maxBlockDimX
+ *      - Expected output: return `hipSuccess`
+ *    -# When blockDimY == maxBlockDimY
+ *      - Expected output: return `hipSuccess`
+ *    -# When blockDimZ == maxBlockDimZ
+ *      - Expected output: return `hipSuccess`
+ *    -# When only start event is passed
+ *      - Expected output: return `hipSucccess`
+ *    -# When only stop event is passed
+ *      - Expected output: return `hipSuccess`
+ * Test source
+ * ------------------------
+ *  - unit/module/hipExtModuleLaunchKernel.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 5.2
+ */
 TEST_CASE("Unit_hipExtModuleLaunchKernel_Positive_Parameters") {
   ModuleLaunchKernelPositiveParameters<hipExtModuleLaunchKernel>();
 
@@ -282,6 +344,56 @@ TEST_CASE("Unit_hipExtModuleLaunchKernel_Positive_Parameters") {
   }
 }
 
+/**
+ * Test Description
+ * ------------------------
+ *  - Validates handling of invalid arguments:
+ *    -# When pointer to the kernel function is `nullptr`
+ *      - Expected output: return `hipErrorInvalidResourceHandle`
+ *    -# When gridDimX == 0
+ *      - Expected output: return `hipErrorInvalidValue`
+ *    -# When gridDimY == 0
+ *      - Expected output: return `hipErrorInvalidValue`
+ *    -# When gridDimZ == 0
+ *      - Expected output: return `hipErrorInvalidValue`
+ *    -# When blockDimX == 0
+ *      - Expected output: return `hipErrorInvalidValue`
+ *    -# When blockDimY == 0
+ *      - Expected output: return `hipErrorInvalidValue`
+ *    -# When blockDimZ == 0
+ *      - Expected output: return `hipErrorInvalidValue`
+ *    -# When gridDimX > maxGridDimX
+ *      - Disabled on AMD because is returns `hipSuccess`
+ *      - Expected output: return `hipErrorInvalidValue`
+ *    -# When gridDimY > maxGridDimY
+ *      - Disabled on AMD because is returns `hipSuccess`
+ *      - Expected output: return `hipErrorInvalidValue`
+ *    -# When gridDimZ > maxGridDimZ
+ *      - Disabled on AMD because is returns `hipSuccess`
+ *      - Expected output: return `hipErrorInvalidValue`
+ *    -# When blockDimX > maxBlockDimX
+ *      - Expected output: return `hipErrorInvalidValue`
+ *    -# When blockDimY > maxBlockDimY
+ *      - Expected output: return `hipErrorInvalidValue`
+ *    -# When blockDimZ > maxBlockDimZ
+ *      - Expected output: return `hipErrorInvalidValue`
+ *    -# When blockDimX * blockDimY * blockDimZ > MaxThreadsPerBlock
+ *      - Expected output: return `hipErrorInvalidValue`
+ *    -# When sharedMemBytes > max shared memory per block
+ *      - Expected output: return `hipErrorInvalidValue`
+ *    -# When stream is not valid, e.g., it is destroyed before kernel launch
+ *      - Expected output: return `hipErrorInvalidValue`
+ *    -# When passing kernelArgs and extra simultaneously
+ *      - Expected output: return `hipErrorInvalidValue`
+ *    -# When extra is not valid
+ *      - Expected output: return `hipErrorInvalidValue`
+ * Test source
+ * ------------------------
+ *  - unit/module/hipExtModuleLaunchKernel.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 5.2
+ */
 TEST_CASE("Unit_hipExtModuleLaunchKernel_Negative_Parameters") {
   ModuleLaunchKernelNegativeParameters<hipExtModuleLaunchKernel>();
 }

catch/unit/module/hipModuleGetFunction.cc

@@ -24,18 +24,61 @@ THE SOFTWARE.
 #include <hip_test_common.hh>
 #include <hip/hip_runtime_api.h>
 
+/**
+ * @addtogroup hipModuleGetFunction hipModuleGetFunction
+ * @{
+ * @ingroup ModuleTest
+ * `hipModuleGetFunction(hipFunction_t* function, hipModule_t module, const char* kname)` -
+ * Function with kname will be extracted if present in module.
+ */
+
 static hipModule_t GetModule() {
   HIP_CHECK(hipFree(nullptr));
   static const auto mg = ModuleGuard::LoadModule("get_function_module.code");
   return mg.module();
 }
 
+/**
+ * Test Description
+ * ------------------------
+ *  - Get function from the loaded module.
+ *  - Checks that the function is not `nullptr`.
+ * Test source
+ * ------------------------
+ *  - unit/module/hipModuleGetFunction.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 5.2
+ */
 TEST_CASE("Unit_hipModuleGetFunction_Positive_Basic") {
   hipFunction_t kernel = nullptr;
   HIP_CHECK(hipModuleGetFunction(&kernel, GetModule(), "GlobalKernel"));
   REQUIRE(kernel != nullptr);
 }
 
+/**
+ * Test Description
+ * ------------------------
+ *  - Verifies handling of invalid arguments:
+ *    -# When output pointer to the function is `nullptr`
+ *      - Expected output: return `hipErrorInvalidValue`
+ *    -# When module pointer is `nullptr`
+ *      - Expected output: return `hipErrorInvalidValue`
+ *    -# When the function name is `nullptr`
+ *      - Expected output: return `hipErrorInvalidValue`
+ *    -# When the function name is an empty string
+ *      - Expected output: return `hipErrorInvalidValue`
+ *    -# When the function name is not existing function
+ *      - Expected output: return `hipErrorNotFound`
+ *    -# When the function name is __device__ function
+ *      - Expected output: return `hipErrorNotFound`
+ * Test source
+ * ------------------------
+ *  - unit/module/hipModuleGetFunction.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 5.2
+ */
 TEST_CASE("Unit_hipModuleGetFunction_Negative_Parameters") {
   hipFunction_t kernel = nullptr;
 
@@ -71,4 +114,4 @@ TEST_CASE("Unit_hipModuleGetFunction_Negative_Parameters") {
   SECTION("kname == __device__ kernel") {
     HIP_CHECK_ERROR(hipModuleGetFunction(&kernel, GetModule(), "DeviceKernel"), hipErrorNotFound);
   }
-}
+}

catch/unit/module/hipModuleGetGlobal.cc

@@ -32,6 +32,16 @@ THE SOFTWARE.
 #include "hip_module_common.hh"
 #include "hipModuleGetGlobal.hh"
 
+/**
+ * @addtogroup hipModuleGetGlobal hipModuleGetGlobal
+ * @{
+ * @ingroup ModuleTest
+ * `hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes, 
+ * hipModule_t hmod, const char* name)` -
+ * Returns a global pointer from a module.
+ * Returns in *dptr and *bytes the pointer and size of the global of name name located in module hmod.
+ */
+
 template <typename T, size_t N>
 static void HipModuleGetGlobalTest(hipModule_t module, const std::string global_name) {
   constexpr auto size = N * sizeof(T);
@@ -82,6 +92,22 @@ static inline hipModule_t GetModule() {
   return mg.module();
 }
 
+/**
+ * Test Description
+ * ------------------------
+ *  - Validates that the API returns correct size and a usable
+ *    pointer to the requested symbol that can be array or scalar:
+ *    -# int
+ *    -# float
+ *    -# char
+ *    -# double
+ * Test source
+ * ------------------------
+ *  - unit/module/hipModuleGetGlobal.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 5.2
+ */
 TEST_CASE("Unit_hipModuleGetGlobal_Positive_Basic") {
   hipModule_t module = GetModule();
 
@@ -94,6 +120,21 @@ TEST_CASE("Unit_hipModuleGetGlobal_Positive_Basic") {
   SECTION("double") { HIP_MODULE_GET_GLOBAL_TEST(double, module); }
 }
 
+/**
+ * Test Description
+ * ------------------------
+ *  - Validates positive handling of `nullptr` arguments:
+ *    -# When the device pointer is `nullptr`
+ *      - Expected output: return `hipSuccess`
+ *    -# When the bytes pointer is `nullptr`
+ *      - Expected output: return `hipSuccess`
+ * Test source
+ * ------------------------
+ *  - unit/module/hipModuleGetGlobal.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 5.2
+ */
 TEST_CASE("Unit_hipModuleGetGlobal_Positive_Parameters") {
   hipModule_t module = GetModule();
   hipDeviceptr_t global = 0;
@@ -108,6 +149,27 @@ TEST_CASE("Unit_hipModuleGetGlobal_Positive_Parameters") {
   }
 }
 
+/**
+ * Test Description
+ * ------------------------
+ *  - Validates handling of invalid arguments:
+ *    -# When device pointer and bytes pointer are `nullptr`
+ *      - Expected output: return `hipErrorInvalidValue`
+ *    -# When module pointer is `nullptr`
+ *      - Expected output: return `hipErrorInvalidResourceHandle`
+ *    -# When module name pointer is `nullptr`
+ *      - Expected output: return `hipErrorInvalidValue`
+ *    -# When module name is an empty string
+ *      - Expected output: return `hipErrorInvalidValue`
+ *    -# When module name is for module that does not exist
+ *      - Expected output: return `hipErrorNotFound`
+ * Test source
+ * ------------------------
+ *  - unit/module/hipModuleGetGlobal.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 5.2
+ */
 TEST_CASE("Unit_hipModuleGetGlobal_Negative_Parameters") {
   hipModule_t module = GetModule();
   hipDeviceptr_t global = 0;
@@ -142,4 +204,4 @@ TEST_CASE("Unit_hipModuleGetGlobal_Negative_Name_Is_Empty_String") {
 TEST_CASE("Unit_hipModuleGetGlobal_Negative_Dptr_And_Bytes_Are_Nullptr") {
   hipModule_t module = GetModule();
   HIP_CHECK_ERROR(hipModuleGetGlobal(nullptr, nullptr, module, "int_var"), hipErrorInvalidValue);
-}
+}