diff --git a/catch/hipTestMain/config/config_amd_linux b/catch/hipTestMain/config/config_amd_linux index 43bbdd1c8..b33784255 100644 --- a/catch/hipTestMain/config/config_amd_linux +++ b/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", diff --git a/catch/hipTestMain/config/config_amd_windows b/catch/hipTestMain/config/config_amd_windows index d7229c192..777ddf3f9 100644 --- a/catch/hipTestMain/config/config_amd_windows +++ b/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", diff --git a/catch/include/hip_test_defgroups.hh b/catch/include/hip_test_defgroups.hh index 680dfa8a0..083556108 100644 --- a/catch/include/hip_test_defgroups.hh +++ b/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 * @{ diff --git a/catch/unit/CMakeLists.txt b/catch/unit/CMakeLists.txt index 304016410..65b009728 100644 --- a/catch/unit/CMakeLists.txt +++ b/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) diff --git a/catch/unit/module/get_tex_ref_module.cc b/catch/unit/module/get_tex_ref_module.cc index edd681ca8..7ea44d757 100644 --- a/catch/unit/module/get_tex_ref_module.cc +++ b/catch/unit/module/get_tex_ref_module.cc @@ -25,4 +25,4 @@ THE SOFTWARE. texture tex; -#endif // CUDA_VERSION < CUDA_12000 \ No newline at end of file +#endif // CUDA_VERSION < CUDA_12000 diff --git a/catch/unit/module/hipExtModuleLaunchKernel.cc b/catch/unit/module/hipExtModuleLaunchKernel.cc index 8c77b796d..876fc39c6 100644 --- a/catch/unit/module/hipExtModuleLaunchKernel.cc +++ b/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(); @@ -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(); @@ -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(); } diff --git a/catch/unit/module/hipModuleGetFunction.cc b/catch/unit/module/hipModuleGetFunction.cc index 676a61c33..1721af0c4 100644 --- a/catch/unit/module/hipModuleGetFunction.cc +++ b/catch/unit/module/hipModuleGetFunction.cc @@ -24,18 +24,61 @@ THE SOFTWARE. #include #include +/** + * @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); } -} \ No newline at end of file +} diff --git a/catch/unit/module/hipModuleGetGlobal.cc b/catch/unit/module/hipModuleGetGlobal.cc index 32f46f0ca..f84c1fb0d 100644 --- a/catch/unit/module/hipModuleGetGlobal.cc +++ b/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 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); -} \ No newline at end of file +} diff --git a/catch/unit/module/hipModuleGetTexRef.cc b/catch/unit/module/hipModuleGetTexRef.cc index 1b3c8a50b..228f83787 100644 --- a/catch/unit/module/hipModuleGetTexRef.cc +++ b/catch/unit/module/hipModuleGetTexRef.cc @@ -24,18 +24,59 @@ THE SOFTWARE. #include #include +/** + * @addtogroup hipModuleGetTexRef hipModuleGetTexRef + * @{ + * @ingroup ModuleTest + * `hipModuleGetTexRef(textureReference** texRef, hipModule_t hmod, const char* name)` - + * Returns the handle of the texture reference with the name from the module. + */ + static hipModule_t GetModule() { HIP_CHECK(hipFree(nullptr)); static const auto mg = ModuleGuard::LoadModule("get_tex_ref_module.code"); return mg.module(); } +/** + * Test Description + * ------------------------ + * - Get the texture reference handle from the loaded module. + * - Validates that texture handle is not `nullptr`. + * Test source + * ------------------------ + * - unit/module/hipModuleGetTexRef.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipModuleGetTexRef_Positive_Basic") { hipTexRef tex_ref = nullptr; HIP_CHECK(hipModuleGetTexRef(&tex_ref, GetModule(), "tex")); REQUIRE(tex_ref != nullptr); } +/** + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# When output pointer to the texture reference is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When module is not loaded and is `nullptr` + * - Expected output: return `hipErrorInvalidResourceHandle` + * -# When texture reference name is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When texture reference name is an empty string + * - Expected output: return `hipErrorInvalidValue` + * -# When texture reference name is the name of non-existing texture reference + * - Expected output: return `hipErrorNotFound` + * Test source + * ------------------------ + * - unit/module/hipModuleGetTexRef.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipModuleGetTexRef_Negative_Parameters") { hipModule_t module = GetModule(); hipTexRef tex_ref = nullptr; @@ -64,4 +105,4 @@ TEST_CASE("Unit_hipModuleGetTexRef_Negative_Name_Is_Empty_String") { hipTexRef tex_ref = nullptr; HIP_CHECK_ERROR(hipModuleGetTexRef(&tex_ref, module, ""), hipErrorInvalidValue); -} \ No newline at end of file +} diff --git a/catch/unit/module/hipModuleLaunchKernel.cc b/catch/unit/module/hipModuleLaunchKernel.cc index f440e8c01..8cfba5abb 100644 --- a/catch/unit/module/hipModuleLaunchKernel.cc +++ b/catch/unit/module/hipModuleLaunchKernel.cc @@ -24,6 +24,17 @@ THE SOFTWARE. #include #include +/** + * @addtogroup hipModuleLaunchKernel hipModuleLaunchKernel + * @{ + * @ingroup ModuleTest + * `hipModuleLaunchKernel(hipFunction_t f, unsigned int gridDimX, unsigned int gridDimY, + * unsigned int gridDimZ, unsigned int blockDimX, unsigned int blockDimY, unsigned int blockDimZ, + * unsigned int sharedMemBytes, hipStream_t stream, void** kernelParams, void** extra)` - + * Launches kernel f with launch parameters and shared memory on stream with arguments passed + * to kernelparams or extra. + */ + static hipError_t hipModuleLaunchKernelWrapper(hipFunction_t f, uint32_t gridX, uint32_t gridY, uint32_t gridZ, uint32_t blockX, uint32_t blockY, uint32_t blockZ, size_t sharedMemBytes, @@ -33,17 +44,107 @@ static hipError_t hipModuleLaunchKernelWrapper(hipFunction_t f, uint32_t gridX, hStream, kernelParams, extra); } +/** + * 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` + * Test source + * ------------------------ + * - unit/module/hipModuleLaunchKernel.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipModuleLaunchKernel_Positive_Basic") { HIP_CHECK(hipFree(nullptr)); ModuleLaunchKernelPositiveBasic(); } +/** + * 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` + * Test source + * ------------------------ + * - unit/module/hipModuleLaunchKernel.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipModuleLaunchKernel_Positive_Parameters") { HIP_CHECK(hipFree(nullptr)); ModuleLaunchKernelPositiveParameters(); } +/** + * 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/hipModuleLaunchKernel.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipModuleLaunchKernel_Negative_Parameters") { HIP_CHECK(hipFree(nullptr)); ModuleLaunchKernelNegativeParameters(); -} \ No newline at end of file +} diff --git a/catch/unit/module/hipModuleLoad.cc b/catch/unit/module/hipModuleLoad.cc index 5812c6a69..79ac7cd10 100644 --- a/catch/unit/module/hipModuleLoad.cc +++ b/catch/unit/module/hipModuleLoad.cc @@ -22,6 +22,26 @@ THE SOFTWARE. #include #include +/** + * @addtogroup hipModuleLoad hipModuleLoad + * @{ + * @ingroup ModuleTest + * `hipModuleLoad(hipModule_t* module, const char* fname)` - + * Loads code object from file into a `hipModule_t`. + */ + +/** + * Test Description + * ------------------------ + * - Loads module from the file. + * - Verifies that the loaded module is not `nullptr`. + * Test source + * ------------------------ + * - unit/module/hipModuleLoad.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipModuleLoad_Positive_Basic") { HIP_CHECK(hipFree(nullptr)); hipModule_t module = nullptr; @@ -30,6 +50,27 @@ TEST_CASE("Unit_hipModuleLoad_Positive_Basic") { HIP_CHECK(hipModuleUnload(module)); } +/** + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# When otput pointer to the module is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When module file name is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When module file name is an empty string + * - Expected output: return `hipErrorInvalidValue` + * -# When module file with the specified name does not exist + * - Expected output: return `hipErrorFileNotFound` + * -# When the file is not a module file + * - Expected output: return `hipErrorInvalidImage` + * Test source + * ------------------------ + * - unit/module/hipModuleLoad.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipModuleLoad_Negative_Parameters") { HIP_CHECK(hipFree(nullptr)); hipModule_t module; @@ -56,4 +97,4 @@ TEST_CASE("Unit_hipModuleLoad_Negative_Load_From_A_File_That_Is_Not_A_Module") { hipModule_t module; HIP_CHECK_ERROR(hipModuleLoad(&module, "not_a_module.txt"), hipErrorInvalidImage); -} \ No newline at end of file +} diff --git a/catch/unit/module/hipModuleLoadData.cc b/catch/unit/module/hipModuleLoadData.cc index 4d364f1d3..fb85bcc7d 100644 --- a/catch/unit/module/hipModuleLoadData.cc +++ b/catch/unit/module/hipModuleLoadData.cc @@ -24,7 +24,32 @@ THE SOFTWARE. #include #include +/** + * @addtogroup hipModuleLoadData hipModuleLoadData + * @{ + * @ingroup ModuleTest + * `hipModuleLoadData(hipModule_t* module, const void* image)` - + * Builds module from code object which resides in host memory. Image is pointer to that + * location. + */ +/** + * Test Description + * ------------------------ + * - Validates different formats of image: + * -# When compiled module is loaded from file + * - Loads module as a binary file into a user space buffer named image. + * - Checks that the module is not `nullptr`. + * -# When module is loaded via RTC + * - Loads RTC module into a user space buffer named image. + * - Checks that the module is not `nullptr` + * Test source + * ------------------------ + * - unit/module/hipModuleLoadData.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipModuleLoadData_Positive_Basic") { HIP_CHECK(hipFree(nullptr)); hipModule_t module = nullptr; @@ -44,6 +69,23 @@ TEST_CASE("Unit_hipModuleLoadData_Positive_Basic") { } } +/** + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# When output pointer to the module is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When pointer to the space buffer named image is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When image is an empty string + * - Expected output: return `hipErrorInvalidValue` + * Test source + * ------------------------ + * - unit/module/hipModuleLoadData.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipModuleLoadData_Negative_Parameters") { HIP_CHECK(hipFree(nullptr)); hipModule_t module; @@ -64,4 +106,4 @@ TEST_CASE("Unit_hipModuleLoadData_Negative_Image_Is_An_Empty_String") { hipModule_t module; HIP_CHECK_ERROR(hipModuleLoadData(&module, ""), hipErrorInvalidImage); -} \ No newline at end of file +} diff --git a/catch/unit/module/hipModuleLoadDataEx.cc b/catch/unit/module/hipModuleLoadDataEx.cc index 2ee3f833f..da5a83c3c 100644 --- a/catch/unit/module/hipModuleLoadDataEx.cc +++ b/catch/unit/module/hipModuleLoadDataEx.cc @@ -24,7 +24,33 @@ THE SOFTWARE. #include #include +/** + * @addtogroup hipModuleLoadDataEx hipModuleLoadDataEx + * @{ + * @ingroup ModuleTest + * `hipModuleLoadDataEx(hipModule_t* module, const void* image, + * unsigned int numOptions, hipJitOption* options, void** optionValues)` - + * Builds module from code object which resides in host memory. Image is pointer to that + * location. Options are not used. hipModuleLoadData is called. + */ +/** + * Test Description + * ------------------------ + * - Validates different formats of image: + * -# When compiled module is loaded from file + * - Loads module as a binary file into a user space buffer named image. + * - Checks that the module is not `nullptr`. + * -# When module is loaded via RTC + * - Loads RTC module into a user space buffer named image. + * - Checks that the module is not `nullptr` + * Test source + * ------------------------ + * - unit/module/hipModuleLoadDataEx.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipModuleLoadDataEx_Positive_Basic") { HIP_CHECK(hipFree(nullptr)); hipModule_t module = nullptr; @@ -44,6 +70,23 @@ TEST_CASE("Unit_hipModuleLoadDataEx_Positive_Basic") { } } +/** + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# When output pointer to the module is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When pointer to the space buffer named image is `nullptr` + * - Expected output: return `hipErrorInvalidValue` + * -# When image is an empty string + * - Expected output: return `hipErrorInvalidValue` + * Test source + * ------------------------ + * - unit/module/hipModuleLoadDataEx.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipModuleLoadDataEx_Negative_Parameters") { HIP_CHECK(hipFree(nullptr)); hipModule_t module = nullptr; @@ -66,4 +109,4 @@ TEST_CASE("Unit_hipModuleLoadDataEx_Negative_Image_Is_An_Empty_String") { hipModule_t module; HIP_CHECK_ERROR(hipModuleLoadDataEx(&module, "", 0, nullptr, nullptr), hipErrorInvalidImage); -} \ No newline at end of file +} diff --git a/catch/unit/module/hipModuleUnload.cc b/catch/unit/module/hipModuleUnload.cc index 914b66c89..87c9b41ae 100644 --- a/catch/unit/module/hipModuleUnload.cc +++ b/catch/unit/module/hipModuleUnload.cc @@ -22,12 +22,50 @@ THE SOFTWARE. #include #include +/** + * @addtogroup hipModuleUnload hipModuleUnload + * @{ + * @ingroup ModuleTest + * `hipModuleUnload(hipModule_t module)` - + * Frees the module. + * ________________________ + * Test cases from other modules: + * - @ref Unit_hipModuleLoad_Positive_Basic + * - @ref Unit_hipModuleLoadData_Positive_Basic + */ + +/** + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# When module is `nullptr` + * - Expected output: return `hipErrorInvalidResourceHandle` + * Test source + * ------------------------ + * - unit/module/hipModuleUnload.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipModuleUnload_Negative_Module_Is_Nullptr") { HIP_CHECK(hipFree(nullptr)); HIP_CHECK_ERROR(hipModuleUnload(nullptr), hipErrorInvalidResourceHandle); } +/** + * Test Description + * ------------------------ + * - Validates handling of invalid arguments: + * -# When module is already unloaded + * - Expected output: return `hipErrorNotFound` + * Test source + * ------------------------ + * - unit/module/hipModuleUnload.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ TEST_CASE("Unit_hipModuleUnload_Negative_Double_Unload") { HIP_CHECK(hipFree(nullptr));