diff --git a/catch/hipTestMain/config/config_amd_linux b/catch/hipTestMain/config/config_amd_linux index 8d891b9e5..5dcbc891d 100644 --- a/catch/hipTestMain/config/config_amd_linux +++ b/catch/hipTestMain/config/config_amd_linux @@ -168,10 +168,6 @@ "Unit_hipGraphicsResourceGetMappedPointer_Negative_Parameters", "Unit_hipGraphicsUnmapResources_Negative_Parameters", "Unit_hipGraphicsUnregisterResource_Negative_Parameters", - "SWDEV-443760: This test fails when device memory is used for kernel args", - "=== Below tests fail in external CI for PR https://github.com/ROCm-Developer-Tools/hip-tests/pull/356 ===", - "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", diff --git a/catch/hipTestMain/config/config_amd_windows b/catch/hipTestMain/config/config_amd_windows index 91a6107dc..5c2d74a27 100644 --- a/catch/hipTestMain/config/config_amd_windows +++ b/catch/hipTestMain/config/config_amd_windows @@ -644,8 +644,6 @@ "Unit___syncthreads_count_Positive_Basic", "Unit___syncthreads_and_Positive_Basic", "Unit___syncthreads_or_Positive_Basic", - "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", diff --git a/catch/hipTestMain/config/config_nvidia_linux.json b/catch/hipTestMain/config/config_nvidia_linux.json index fc8a35dd1..c3827978f 100644 --- a/catch/hipTestMain/config/config_nvidia_linux.json +++ b/catch/hipTestMain/config/config_nvidia_linux.json @@ -70,12 +70,6 @@ "Unit_hipMemcpy2D_Positive_Synchronization_Behavior", "Unit_hipDrvMemcpy3D_Positive_Synchronization_Behavior", "Unit_hipFreeMipmappedArray_Negative_DoubleFree", - "Unit_hipModuleLoad_Positive_Basic", - "Unit_hipModuleLoad_Negative_Load_From_A_File_That_Is_Not_A_Module", - "Unit_hipModuleLoadData_Positive_Basic", - "Unit_hipModuleLoadData_Negative_Parameters", - "Unit_hipModuleLoadDataEx_Positive_Basic", - "Unit_hipModuleLoadDataEx_Negative_Parameters", "Performance_hipMemsetD16", "Performance_hipMemsetD16Async", "Performance_hipMemsetD32", diff --git a/catch/unit/module/CMakeLists.txt b/catch/unit/module/CMakeLists.txt index 3f891d974..4720ae3bc 100644 --- a/catch/unit/module/CMakeLists.txt +++ b/catch/unit/module/CMakeLists.txt @@ -36,185 +36,96 @@ set(TEST_SRC hipDrvLaunchKernelEx.cc ) -add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/get_function_module.code - COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} --std=c++17 ${CMAKE_CURRENT_SOURCE_DIR}/get_function_module.cc - -o get_function_module.code - -I${HIP_PATH}/include/ --hip-path=${HIP_PATH} - DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/get_function_module.cc) -add_custom_target(get_function_module ALL DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/get_function_module.code) - -add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/launch_kernel_module.code - COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} --std=c++17 ${CMAKE_CURRENT_SOURCE_DIR}/launch_kernel_module.cc - -o launch_kernel_module.code - -I${HIP_PATH}/include/ --hip-path=${HIP_PATH} - DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/launch_kernel_module.cc) -add_custom_target(launch_kernel_module ALL DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/launch_kernel_module.code) - -add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/get_global_test_module.code - COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} --std=c++17 ${CMAKE_CURRENT_SOURCE_DIR}/get_global_test_module.cc - -o get_global_test_module.code - -I${HIP_PATH}/include/ --hip-path=${HIP_PATH} - DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/get_global_test_module.cc) -add_custom_target(get_global_test_module ALL DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/get_global_test_module.code) - -add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/get_tex_ref_module.code - COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} --std=c++17 ${CMAKE_CURRENT_SOURCE_DIR}/get_tex_ref_module.cc - -o get_tex_ref_module.code - -I${HIP_PATH}/include/ --hip-path=${HIP_PATH} - DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/get_tex_ref_module.cc) -add_custom_target(get_tex_ref_module ALL DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/get_tex_ref_module.code) - -add_custom_target(coopKernel.code - COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} - ${CMAKE_CURRENT_SOURCE_DIR}/coopKernel.cpp - -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/module/coopKernel.code - -I${HIP_PATH}/include/ --rocm-path=${ROCM_PATH} - -I${CMAKE_CURRENT_SOURCE_DIR}/../../include) - -set_property(GLOBAL APPEND PROPERTY G_INSTALL_CUSTOM_TARGETS - ${CMAKE_CURRENT_BINARY_DIR}/get_function_module.code - ${CMAKE_CURRENT_BINARY_DIR}/launch_kernel_module.code - ${CMAKE_CURRENT_BINARY_DIR}/get_global_test_module.code - ${CMAKE_CURRENT_BINARY_DIR}/get_tex_ref_module.code - ${CMAKE_CURRENT_BINARY_DIR}/coopKernel.code -) -# Note to pass arch use format like -DOFFLOAD_ARCH_STR="--offload-arch=gfx900 --offload-arch=gfx906" -# having space at the start/end of OFFLOAD_ARCH_STR can cause build failures +function(add_custom_module target module_name source_file_name flags type) + set(output_file ${CMAKE_CURRENT_BINARY_DIR}/${module_name}.${type}) + set(source_file ${CMAKE_CURRENT_SOURCE_DIR}/${source_file_name}) + + add_custom_command( + OUTPUT ${output_file} + COMMAND ${CMAKE_CXX_COMPILER} ${flags} ${source_file} -o ${output_file} + -I${HIP_PATH}/include/ -I${CMAKE_CURRENT_SOURCE_DIR}/../../include --hip-path=${HIP_PATH} + DEPENDS ${source_file} + ) + + set(custom_target_name ${module_name}_${type}) + add_custom_target(${custom_target_name} ALL DEPENDS ${output_file}) + set_property(GLOBAL APPEND PROPERTY G_INSTALL_CUSTOM_TARGETS ${output_file}) + add_dependencies(${target} ${custom_target_name}) +endfunction() + +function(add_custom_file target module_name source_file_name flags) + set(output_file ${CMAKE_CURRENT_BINARY_DIR}/${module_name}.txt) + set(source_file ${CMAKE_CURRENT_SOURCE_DIR}/${source_file_name}) + + add_custom_command( + OUTPUT ${output_file} + COMMAND ${CMAKE_COMMAND} -E copy ${source_file} ${output_file} + DEPENDS ${source_file} + ) + + add_custom_target(${module_name} ALL DEPENDS ${output_file}) + set_property(GLOBAL APPEND PROPERTY G_INSTALL_CUSTOM_TARGETS ${output_file}) + add_dependencies(${target} ${module_name}) +endfunction() + +add_custom_module(build_tests get_function_module get_function_module.cc "--genco;${OFFLOAD_ARCH_STR}" "code") +add_custom_module(build_tests launch_kernel_module launch_kernel_module.cc "--genco;${OFFLOAD_ARCH_STR}" "code") +add_custom_module(build_tests get_global_test_module get_global_test_module.cc "--genco;${OFFLOAD_ARCH_STR}" "code") +add_custom_module(build_tests get_tex_ref_module get_tex_ref_module.cc "--genco;${OFFLOAD_ARCH_STR}" "code") +add_custom_module(build_tests empty_module empty_module.cc "--genco;${OFFLOAD_ARCH_STR}" "code") +add_custom_module(build_tests coopKernel coopKernel.cpp "--genco;${OFFLOAD_ARCH_STR}" "code") + +add_custom_file(build_tests empty_file empty_file.txt "") +add_custom_file(build_tests not_a_module not_a_module.txt "") if(HIP_PLATFORM MATCHES "amd") -set(TEST_SRC - ${TEST_SRC} - hipExtModuleLaunchKernel.cc - hipHccModuleLaunchKernel.cc - hipGetProcAddressModuleApis.cc) - -add_custom_target(empty_module.code - COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} - ${CMAKE_CURRENT_SOURCE_DIR}/empty_module.cc - -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/module/empty_module.code - -I${HIP_PATH}/include/ --hip-path=${HIP_PATH} - -I${CMAKE_CURRENT_SOURCE_DIR}/../../include) - -add_custom_target(copyKernel.code - COMMAND ${CMAKE_CXX_COMPILER} -mcode-object-version=5 --genco ${OFFLOAD_ARCH_STR} - ${CMAKE_CURRENT_SOURCE_DIR}/copyKernel.cc - -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/module/copyKernel.code - -I${HIP_PATH}/include/ --hip-path=${HIP_PATH} - -I${CMAKE_CURRENT_SOURCE_DIR}/../../include) - -add_custom_target(copyKernel.s - COMMAND ${CMAKE_CXX_COMPILER} -mcode-object-version=5 -S ${CMAKE_CURRENT_SOURCE_DIR}/copyKernel.cc - -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/module/copyKernel.s - -I${HIP_PATH}/include/ --hip-path=${HIP_PATH} - -I${CMAKE_CURRENT_SOURCE_DIR}/../../include) - -add_custom_target(addKernel.code - COMMAND ${CMAKE_CXX_COMPILER} -mcode-object-version=5 --genco ${OFFLOAD_ARCH_STR} - ${CMAKE_CURRENT_SOURCE_DIR}/addKernel.cc - -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/module/addKernel.code - -I${HIP_PATH}/include/ --hip-path=${HIP_PATH} - -I${CMAKE_CURRENT_SOURCE_DIR}/../../include) - -add_custom_target(copyKernelCompressed.code - COMMAND ${CMAKE_CXX_COMPILER} -mcode-object-version=5 --offload-compress --genco ${OFFLOAD_ARCH_STR} - ${CMAKE_CURRENT_SOURCE_DIR}/copyKernel.cc - -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/module/copyKernelCompressed.code - -I${HIP_PATH}/include/ --hip-path=${HIP_PATH} - -I${CMAKE_CURRENT_SOURCE_DIR}/../../include) - -set(OFFLOAD_ARCH_GENERIC_STR "--offload-arch=gfx9-generic --offload-arch=gfx9-4-generic:sramecc+:xnack- --offload-arch=gfx9-4-generic:sramecc-:xnack- --offload-arch=gfx9-4-generic:xnack+ --offload-arch=gfx10-1-generic --offload-arch=gfx10-3-generic --offload-arch=gfx11-generic --offload-arch=gfx12-generic") -add_custom_target(copyKernelGenericTarget.code - COMMAND ${CMAKE_CXX_COMPILER} -mcode-object-version=6 --genco ${OFFLOAD_ARCH_GENERIC_STR} - ${CMAKE_CURRENT_SOURCE_DIR}/copyKernel.cc - -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/module/copyKernelGenericTarget.code - -I${HIP_PATH}/include/ --hip-path=${HIP_PATH} - -I${CMAKE_CURRENT_SOURCE_DIR}/../../include) - -add_custom_target(copyKernelGenericTargetCompressed.code - COMMAND ${CMAKE_CXX_COMPILER} -mcode-object-version=6 --offload-compress --genco ${OFFLOAD_ARCH_GENERIC_STR} - ${CMAKE_CURRENT_SOURCE_DIR}/copyKernel.cc - -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/module/copyKernelGenericTargetCompressed.code - -I${HIP_PATH}/include/ --hip-path=${HIP_PATH} - -I${CMAKE_CURRENT_SOURCE_DIR}/../../include) -set_property(GLOBAL APPEND PROPERTY G_INSTALL_CUSTOM_TARGETS - ${CMAKE_CURRENT_BINARY_DIR}/empty_module.code - ${CMAKE_CURRENT_BINARY_DIR}/copyKernel.code - ${CMAKE_CURRENT_BINARY_DIR}/copyKernel.s - ${CMAKE_CURRENT_BINARY_DIR}/addKernel.code - ${CMAKE_CURRENT_BINARY_DIR}/copyKernelCompressed.code - ${CMAKE_CURRENT_BINARY_DIR}/copyKernelGenericTarget.code - ${CMAKE_CURRENT_BINARY_DIR}/copyKernelGenericTargetCompressed.code -) - -if(UNIX) -set(TEST_SRC - ${TEST_SRC} - hipKerArgOptimization.cc) - -add_custom_target(copiousArgKernel.code - COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} - ${CMAKE_CURRENT_SOURCE_DIR}/copiousArgKernel.cc - -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/module/copiousArgKernel.code - -I${HIP_PATH}/include/ --rocm-path=${ROCM_PATH} - -I${CMAKE_CURRENT_SOURCE_DIR}/../../include --rocm-path=${ROCM_PATH}) - -add_custom_target(copiousArgKernel0.code - COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} - -mllvm -amdgpu-kernarg-preload-count=0 - ${CMAKE_CURRENT_SOURCE_DIR}/copiousArgKernel.cc - -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/module/copiousArgKernel0.code - -I${HIP_PATH}/include/ --rocm-path=${ROCM_PATH} - -I${CMAKE_CURRENT_SOURCE_DIR}/../../include --rocm-path=${ROCM_PATH}) - -add_custom_target(copiousArgKernel1.code - COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} - -mllvm -amdgpu-kernarg-preload-count=1 - ${CMAKE_CURRENT_SOURCE_DIR}/copiousArgKernel.cc - -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/module/copiousArgKernel1.code - -I${HIP_PATH}/include/ --rocm-path=${ROCM_PATH} - -I${CMAKE_CURRENT_SOURCE_DIR}/../../include --rocm-path=${ROCM_PATH}) - -add_custom_target(copiousArgKernel2.code - COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} - -mllvm -amdgpu-kernarg-preload-count=2 - ${CMAKE_CURRENT_SOURCE_DIR}/copiousArgKernel.cc - -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/module/copiousArgKernel2.code - -I${HIP_PATH}/include/ --rocm-path=${ROCM_PATH} - -I${CMAKE_CURRENT_SOURCE_DIR}/../../include --rocm-path=${ROCM_PATH}) - -add_custom_target(copiousArgKernel3.code - COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} - -mllvm -amdgpu-kernarg-preload-count=3 - ${CMAKE_CURRENT_SOURCE_DIR}/copiousArgKernel.cc - -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/module/copiousArgKernel3.code - -I${HIP_PATH}/include/ --rocm-path=${ROCM_PATH} - -I${CMAKE_CURRENT_SOURCE_DIR}/../../include --rocm-path=${ROCM_PATH}) - -add_custom_target(copiousArgKernel16.code - COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} - -mllvm -amdgpu-kernarg-preload-count=16 - ${CMAKE_CURRENT_SOURCE_DIR}/copiousArgKernel.cc - -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/module/copiousArgKernel16.code - -I${HIP_PATH}/include/ --rocm-path=${ROCM_PATH} - -I${CMAKE_CURRENT_SOURCE_DIR}/../../include --rocm-path=${ROCM_PATH}) - -add_custom_target(copiousArgKernel17.code - COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} - -mllvm -amdgpu-kernarg-preload-count=17 - ${CMAKE_CURRENT_SOURCE_DIR}/copiousArgKernel.cc - -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/module/copiousArgKernel17.code - -I${HIP_PATH}/include/ --rocm-path=${ROCM_PATH} - -I${CMAKE_CURRENT_SOURCE_DIR}/../../include --rocm-path=${ROCM_PATH}) -set_property(GLOBAL APPEND PROPERTY G_INSTALL_CUSTOM_TARGETS - ${CMAKE_CURRENT_BINARY_DIR}/copiousArgKernel.code - ${CMAKE_CURRENT_BINARY_DIR}/copiousArgKernel0.code - ${CMAKE_CURRENT_BINARY_DIR}/copiousArgKernel1.code - ${CMAKE_CURRENT_BINARY_DIR}/copiousArgKernel2.code - ${CMAKE_CURRENT_BINARY_DIR}/copiousArgKernel3.code - ${CMAKE_CURRENT_BINARY_DIR}/copiousArgKernel16.code - ${CMAKE_CURRENT_BINARY_DIR}/copiousArgKernel17.code -) -endif() + set(TEST_SRC ${TEST_SRC} hipExtModuleLaunchKernel.cc + hipHccModuleLaunchKernel.cc hipGetProcAddressModuleApis.cc) + + set(OFFLOAD_ARCH_GENERIC_STR + "--offload-arch=gfx9-generic --offload-arch=gfx9-4-generic:sramecc+:xnack- \ + --offload-arch=gfx9-4-generic:sramecc-:xnack- --offload-arch=gfx9-4-generic:xnack+ \ + --offload-arch=gfx10-1-generic --offload-arch=gfx10-3-generic --offload-arch=gfx11-generic \ + --offload-arch=gfx12-generic") + + add_custom_module(build_tests addKernel addKernel.cc + "-mcode-object-version=5;--genco;${OFFLOAD_ARCH_STR};" "code") + add_custom_module(build_tests copyKernelCompressed copyKernel.cc + "-mcode-object-version=5;--genco;--offload-compress;${OFFLOAD_ARCH_STR};" "code") + add_custom_module(build_tests copyKernel copyKernel.cc + "-mcode-object-version=5;--genco;${OFFLOAD_ARCH_STR};" "code") + add_custom_module(build_tests copyKernel copyKernel.cc + "-mcode-object-version=5;-S;" "s") + add_custom_module(build_tests copyKernelGenericTarget copyKernel.cc + "-mcode-object-version=6;--genco;${OFFLOAD_ARCH_GENERIC_STR};" "code") + add_custom_module(build_tests copyKernelGenericTargetCompressed copyKernel.cc + "-mcode-object-version=6;--offload-compress;--genco;${OFFLOAD_ARCH_GENERIC_STR}" "code") + + if(UNIX) + set(TEST_SRC ${TEST_SRC} hipKerArgOptimization.cc) + + add_custom_module(build_tests copiousArgKernel copiousArgKernel.cc + "--genco;${OFFLOAD_ARCH_STR}" "code") + add_custom_module(build_tests copiousArgKernel0 copiousArgKernel.cc + "-mllvm;-amdgpu-kernarg-preload-count=0;--genco;${OFFLOAD_ARCH_STR}" "code" + ) + add_custom_module( + build_tests copiousArgKernel1 copiousArgKernel.cc + "-mllvm;-amdgpu-kernarg-preload-count=1;--genco;${OFFLOAD_ARCH_STR}" "code" + ) + add_custom_module(build_tests copiousArgKernel2 copiousArgKernel.cc + "-mllvm;-amdgpu-kernarg-preload-count=2;--genco;${OFFLOAD_ARCH_STR}" "code" + ) + add_custom_module(build_tests copiousArgKernel3 copiousArgKernel.cc + "-mllvm;-amdgpu-kernarg-preload-count=3;--genco;${OFFLOAD_ARCH_STR}" "code" + ) + add_custom_module(build_tests copiousArgKernel16 copiousArgKernel.cc + "-mllvm;-amdgpu-kernarg-preload-count=16;--genco;${OFFLOAD_ARCH_STR}" "code" + ) + add_custom_module(build_tests copiousArgKernel17 copiousArgKernel.cc + "-mllvm;-amdgpu-kernarg-preload-count=17;--genco;${OFFLOAD_ARCH_STR}" "code" + ) + endif() endif() if(HIP_PLATFORM MATCHES "amd") @@ -230,26 +141,6 @@ hip_add_exe_to_target(NAME ModuleTest COMMON_SHARED_SRC ${COMMON_SHARED_SRC} COMPILE_OPTIONS -std=c++17) -add_dependencies(ModuleTest coopKernel.code) -add_dependencies(ModuleTest get_function_module) -add_dependencies(ModuleTest launch_kernel_module) -add_dependencies(ModuleTest get_global_test_module) -add_dependencies(ModuleTest get_tex_ref_module) - -if(HIP_PLATFORM MATCHES "amd") -add_dependencies(build_tests empty_module.code) -add_dependencies(build_tests copyKernel.code copyKernel.s) -add_dependencies(build_tests addKernel.code) -add_dependencies(build_tests copyKernelCompressed.code) -add_dependencies(build_tests copyKernelGenericTarget.code) -add_dependencies(build_tests copyKernelGenericTargetCompressed.code) - -if(UNIX) -add_dependencies(build_tests copiousArgKernel.code copiousArgKernel0.code copiousArgKernel1.code copiousArgKernel2.code -copiousArgKernel3.code copiousArgKernel16.code copiousArgKernel17.code) -endif() -endif() - add_executable(hipGetFuncBySymbol_exe EXCLUDE_FROM_ALL hipGetFuncBySymbol_exe.cc) add_dependencies(build_tests hipGetFuncBySymbol_exe) set_property(GLOBAL APPEND PROPERTY G_INSTALL_EXE_TARGETS hipGetFuncBySymbol_exe) @@ -276,34 +167,26 @@ hip_add_exe_to_target(NAME module TEST_SRC ${TEST_SRC} TEST_TARGET_NAME build_tests) -add_custom_target(managed_kernel.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} ${CMAKE_CURRENT_SOURCE_DIR}/managed_kernel.cpp -o ${CMAKE_CURRENT_BINARY_DIR}/../module/managed_kernel.code -I${HIP_PATH}/include -I${CMAKE_CURRENT_SOURCE_DIR}/../../include --hip-path=${HIP_PATH}) - hip_add_exe_to_target(NAME managedKernel TEST_SRC ${LINUX_TEST_SRC} TEST_TARGET_NAME build_tests LINKER_LIBS ${CMAKE_DL_LIBS}) -add_custom_target(vcpy_kernel.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} ${CMAKE_CURRENT_SOURCE_DIR}/vcpy_kernel.cpp -o ${CMAKE_CURRENT_BINARY_DIR}/../module/vcpy_kernel.code -I${HIP_PATH}/include -I${CMAKE_CURRENT_SOURCE_DIR}/../../include --hip-path=${HIP_PATH}) - hip_add_exe_to_target(NAME VcpyKernel TEST_SRC ${LINUX_TEST_SRC} TEST_TARGET_NAME build_tests LINKER_LIBS ${CMAKE_DL_LIBS}) -add_custom_target(matmul.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} ${CMAKE_CURRENT_SOURCE_DIR}/matmul.cpp -o ${CMAKE_CURRENT_BINARY_DIR}/../module/matmul.code -I${HIP_PATH}/include -I${CMAKE_CURRENT_SOURCE_DIR}/../../include --hip-path=${HIP_PATH}) - hip_add_exe_to_target(NAME matmul TEST_SRC ${LINUX_TEST_SRC} TEST_TARGET_NAME build_tests LINKER_LIBS ${CMAKE_DL_LIBS}) -add_custom_target(kernel_composite_test.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} ${CMAKE_CURRENT_SOURCE_DIR}/kernel_composite_test.cpp -o ${CMAKE_CURRENT_BINARY_DIR}/../module/kernel_composite_test.code -I${HIP_PATH}/include -I${CMAKE_CURRENT_SOURCE_DIR}/../../include --hip-path=${HIP_PATH}) -set_property(GLOBAL APPEND PROPERTY G_INSTALL_CUSTOM_TARGETS - ${CMAKE_CURRENT_BINARY_DIR}/kernel_composite_test.code - ${CMAKE_CURRENT_BINARY_DIR}/matmul.code - ${CMAKE_CURRENT_BINARY_DIR}/vcpy_kernel.code - ${CMAKE_CURRENT_BINARY_DIR}/managed_kernel.code -) +add_custom_module(module managed_kernel managed_kernel.cpp "--genco;${OFFLOAD_ARCH_STR}" "code") +add_custom_module(module vcpy_kernel vcpy_kernel.cpp "--genco;${OFFLOAD_ARCH_STR}" "code") +add_custom_module(module matmul matmul.cpp "--genco;${OFFLOAD_ARCH_STR}" "code") +add_custom_module(module kernel_composite_test kernel_composite_test.cpp "--genco;${OFFLOAD_ARCH_STR}" "code") + add_executable(testhipModuleLoadUnloadFunc_exe EXCLUDE_FROM_ALL testhipModuleLoadUnloadFunc_exe.cc) set_property(GLOBAL APPEND PROPERTY G_INSTALL_EXE_TARGETS testhipModuleLoadUnloadFunc_exe) @@ -311,4 +194,5 @@ hip_add_exe_to_target(NAME compositeKernel TEST_SRC ${LINUX_TEST_SRC} TEST_TARGET_NAME build_tests LINKER_LIBS ${CMAKE_DL_LIBS}) -add_dependencies(module managed_kernel.code vcpy_kernel.code matmul.code kernel_composite_test.code testhipModuleLoadUnloadFunc_exe) + +add_dependencies(module testhipModuleLoadUnloadFunc_exe) diff --git a/catch/unit/module/empty_file.txt b/catch/unit/module/empty_file.txt new file mode 100644 index 000000000..e69de29bb diff --git a/catch/unit/module/hipExtModuleLaunchKernel.cc b/catch/unit/module/hipExtModuleLaunchKernel.cc index 4faf712d9..221a169c4 100644 --- a/catch/unit/module/hipExtModuleLaunchKernel.cc +++ b/catch/unit/module/hipExtModuleLaunchKernel.cc @@ -265,7 +265,7 @@ TEST_CASE("Unit_hipExtModuleLaunchKernel_Positive_Parameters") { SECTION("Pass only start event") { hipEvent_t start_event = nullptr; HIP_CHECK(hipEventCreate(&start_event)); - const auto kernel = GetKernel(mg.module(), "NOPKernel"); + const auto kernel = GetKernel(GetModule(), "NOPKernel"); HIP_CHECK(hipExtModuleLaunchKernel(kernel, 1, 1, 1, 1, 1, 1, 0, nullptr, nullptr, nullptr, start_event, nullptr)); @@ -277,7 +277,7 @@ TEST_CASE("Unit_hipExtModuleLaunchKernel_Positive_Parameters") { SECTION("Pass only stop event") { hipEvent_t stop_event = nullptr; HIP_CHECK(hipEventCreate(&stop_event)); - const auto kernel = GetKernel(mg.module(), "NOPKernel"); + const auto kernel = GetKernel(GetModule(), "NOPKernel"); HIP_CHECK(hipExtModuleLaunchKernel(kernel, 1, 1, 1, 1, 1, 1, 0, nullptr, nullptr, nullptr, nullptr, stop_event)); diff --git a/catch/unit/module/hipModuleGetGlobal.cc b/catch/unit/module/hipModuleGetGlobal.cc index a4a5423d1..d7cd23149 100644 --- a/catch/unit/module/hipModuleGetGlobal.cc +++ b/catch/unit/module/hipModuleGetGlobal.cc @@ -127,8 +127,10 @@ TEST_CASE("Unit_hipModuleGetGlobal_Negative_Hmod_Is_Nullptr") { hipDeviceptr_t global = 0; size_t global_size = 0; + CTX_CREATE(); HIP_CHECK_ERROR(hipModuleGetGlobal(&global, &global_size, nullptr, "int_var"), hipErrorInvalidResourceHandle); + CTX_DESTROY(); } TEST_CASE("Unit_hipModuleGetGlobal_Negative_Name_Is_Empty_String") { diff --git a/catch/unit/module/hipModuleGetTexRef.cc b/catch/unit/module/hipModuleGetTexRef.cc index 651856bf3..d835e27a5 100644 --- a/catch/unit/module/hipModuleGetTexRef.cc +++ b/catch/unit/module/hipModuleGetTexRef.cc @@ -59,7 +59,9 @@ TEST_CASE("Unit_hipModuleGetTexRef_Negative_Hmod_Is_Nullptr") { CHECK_IMAGE_SUPPORT hipTexRef tex_ref = nullptr; + CTX_CREATE(); HIP_CHECK_ERROR(hipModuleGetTexRef(&tex_ref, nullptr, "tex"), hipErrorInvalidResourceHandle); + CTX_DESTROY(); } TEST_CASE("Unit_hipModuleGetTexRef_Negative_Name_Is_Empty_String") { @@ -68,4 +70,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/hipModuleLaunchCooperativeKernel.cc b/catch/unit/module/hipModuleLaunchCooperativeKernel.cc index 92924d88a..a636bc72c 100644 --- a/catch/unit/module/hipModuleLaunchCooperativeKernel.cc +++ b/catch/unit/module/hipModuleLaunchCooperativeKernel.cc @@ -56,13 +56,13 @@ TEST_CASE("Unit_hipModuleLaunchCooperativeKernel_Positive_Basic") { } SECTION("Cooperative kernel with no arguments") { - hipFunction_t f = GetKernel(mg.module(), "CoopKernel"); + hipFunction_t f = GetKernel(GetModule(), "CoopKernel"); HIP_CHECK(hipModuleLaunchCooperativeKernel(f, 2, 2, 1, 1, 1, 1, 0, nullptr, nullptr)); HIP_CHECK(hipDeviceSynchronize()); } SECTION("Kernel with arguments using kernelParams") { - hipFunction_t f = GetKernel(mg.module(), "Kernel42"); + hipFunction_t f = GetKernel(GetModule(), "Kernel42"); LinearAllocGuard result_dev(LinearAllocs::hipMalloc, sizeof(int)); HIP_CHECK(hipMemset(result_dev.ptr(), 0, sizeof(*result_dev.ptr()))); @@ -94,7 +94,7 @@ TEST_CASE("Unit_hipModuleLaunchCooperativeKernel_Positive_Parameters") { return; } - hipFunction_t f = GetKernel(mg.module(), "NOPKernel"); + hipFunction_t f = GetKernel(GetModule(), "NOPKernel"); SECTION("blockDim.x == maxBlockDimX") { const unsigned int x = GetDeviceAttribute(hipDeviceAttributeMaxBlockDimX, 0); @@ -129,7 +129,7 @@ TEST_CASE("Unit_hipModuleLaunchCooperativeKernel_Negative_Parameters") { return; } - hipFunction_t f = GetKernel(mg.module(), "NOPKernel"); + hipFunction_t f = GetKernel(GetModule(), "NOPKernel"); SECTION("f == nullptr") { HIP_CHECK_ERROR( diff --git a/catch/unit/module/hipModuleLaunchCooperativeKernelMultiDevice.cc b/catch/unit/module/hipModuleLaunchCooperativeKernelMultiDevice.cc index e9f2389ce..be3c52d4f 100644 --- a/catch/unit/module/hipModuleLaunchCooperativeKernelMultiDevice.cc +++ b/catch/unit/module/hipModuleLaunchCooperativeKernelMultiDevice.cc @@ -163,7 +163,7 @@ TEST_CASE("Unit_hipModuleLaunchCooperativeKernelMultiDevice_Negative_Parameters" if (device_count > 1) { SECTION("launchParamsList.func doesn't match across all devices") { - params_list[1].function = GetKernel(mg.module(), "NOPKernel"); + params_list[1].function = GetKernel(GetModule(), "NOPKernel"); #if HT_AMD HIP_CHECK_ERROR( hipModuleLaunchCooperativeKernelMultiDevice(params_list.data(), device_count, 0u), diff --git a/catch/unit/module/hipModuleLoad.cc b/catch/unit/module/hipModuleLoad.cc index 5812c6a69..ef197ba5b 100644 --- a/catch/unit/module/hipModuleLoad.cc +++ b/catch/unit/module/hipModuleLoad.cc @@ -56,4 +56,5 @@ 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 + HIP_CHECK_ERROR(hipModuleLoad(&module, "empty_file.txt"), hipErrorInvalidImage); +} diff --git a/catch/unit/module/hipModuleUnload.cc b/catch/unit/module/hipModuleUnload.cc index 54c1a46cd..a96d2e0a9 100644 --- a/catch/unit/module/hipModuleUnload.cc +++ b/catch/unit/module/hipModuleUnload.cc @@ -32,7 +32,11 @@ TEST_CASE("Unit_hipModuleUnload_Negative_Double_Unload") { hipModule_t module = nullptr; HIP_CHECK(hipModuleLoad(&module, "empty_module.code")); HIP_CHECK(hipModuleUnload(module)); +#if HT_AMD HIP_CHECK_ERROR(hipModuleUnload(module), hipErrorNotFound); +#else + HIP_CHECK_ERROR(hipModuleUnload(module), hipErrorInvalidResourceHandle); +#endif } /** * @addtogroup hipModuleUnload @@ -54,9 +58,11 @@ TEST_CASE("Unit_hipModuleUnload_Negative_Double_Unload") { * - HIP_VERSION >= 5.6 */ TEST_CASE("Unit_hipModuleLoad_basic") { - constexpr auto fileName = "vcpy_kernel.code"; - hipModule_t module; - HIP_CHECK(hipModuleLoad(&module, fileName)); - REQUIRE(module != nullptr); - HIP_CHECK(hipModuleUnload(module)); + CTX_CREATE(); + constexpr auto fileName = "vcpy_kernel.code"; + hipModule_t module; + HIP_CHECK(hipModuleLoad(&module, fileName)); + REQUIRE(module != nullptr); + HIP_CHECK(hipModuleUnload(module)); + CTX_DESTROY(); } diff --git a/catch/unit/module/hip_module_common.cc b/catch/unit/module/hip_module_common.cc index 2cdb250b3..fbb831900 100644 --- a/catch/unit/module/hip_module_common.cc +++ b/catch/unit/module/hip_module_common.cc @@ -49,12 +49,12 @@ ModuleGuard ModuleGuard::LoadModuleDataRTC(const char* code) { // Load module into buffer instead of mapping file to avoid platform specific mechanisms std::vector LoadModuleIntoBuffer(const char* path_string) { - fs::path p(path_string); - const auto file_size = fs::file_size(p); - std::ifstream f(p, std::ios::binary | std::ios::in); - REQUIRE(f); - std::vector empty_module(file_size); - REQUIRE(f.read(empty_module.data(), file_size)); + std::ifstream file_stream(path_string, std::ios::binary | std::ios::in); + REQUIRE(file_stream); + std::vector empty_module((std::istreambuf_iterator(file_stream)), + std::istreambuf_iterator()); + file_stream.close(); + empty_module.push_back('\0'); return empty_module; } @@ -68,4 +68,4 @@ std::vector CreateRTCCharArray(const char* src) { HIPRTC_CHECK(hiprtcGetCode(prog, code.data())); HIPRTC_CHECK(hiprtcDestroyProgram(&prog)); return code; -} \ No newline at end of file +} diff --git a/catch/unit/module/hip_module_launch_kernel_common.hh b/catch/unit/module/hip_module_launch_kernel_common.hh index b5dc9f772..56f40e295 100644 --- a/catch/unit/module/hip_module_launch_kernel_common.hh +++ b/catch/unit/module/hip_module_launch_kernel_common.hh @@ -27,26 +27,25 @@ THE SOFTWARE. #include #include -inline ModuleGuard InitModule() { +static hipModule_t GetModule() { HIP_CHECK(hipFree(nullptr)); - return ModuleGuard::LoadModule("launch_kernel_module.code"); + static const auto mg = ModuleGuard::LoadModule("launch_kernel_module.code"); + return mg.module(); } -inline ModuleGuard mg{InitModule()}; - using ExtModuleLaunchKernelSig = hipError_t(hipFunction_t, uint32_t, uint32_t, uint32_t, uint32_t, uint32_t, uint32_t, size_t, hipStream_t, void**, void**, hipEvent_t, hipEvent_t, uint32_t); template void ModuleLaunchKernelPositiveBasic() { SECTION("Kernel with no arguments") { - hipFunction_t f = GetKernel(mg.module(), "NOPKernel"); + hipFunction_t f = GetKernel(GetModule(), "NOPKernel"); HIP_CHECK(func(f, 1, 1, 1, 1, 1, 1, 0, nullptr, nullptr, nullptr, nullptr, nullptr, 0u)); HIP_CHECK(hipDeviceSynchronize()); } SECTION("Kernel with arguments using kernelParams") { - hipFunction_t f = GetKernel(mg.module(), "Kernel42"); + hipFunction_t f = GetKernel(GetModule(), "Kernel42"); LinearAllocGuard result_dev(LinearAllocs::hipMalloc, sizeof(int)); HIP_CHECK(hipMemset(result_dev.ptr(), 0, sizeof(*result_dev.ptr()))); int* result_ptr = result_dev.ptr(); @@ -58,7 +57,7 @@ template void ModuleLaunchKernelPositiveBasic() } SECTION("Kernel with arguments using extra") { - hipFunction_t f = GetKernel(mg.module(), "Kernel42"); + hipFunction_t f = GetKernel(GetModule(), "Kernel42"); LinearAllocGuard result_dev(LinearAllocs::hipMalloc, sizeof(int)); HIP_CHECK(hipMemset(result_dev.ptr(), 0, sizeof(*result_dev.ptr()))); int* result_ptr = result_dev.ptr(); @@ -81,7 +80,7 @@ template void ModuleLaunchKernelPositiveParamet const auto LaunchNOPKernel = [=](unsigned int gridDimX, unsigned int gridDimY, unsigned int gridDimZ, unsigned int blockDimX, unsigned int blockDimY, unsigned int blockDimZ) { - hipFunction_t f = GetKernel(mg.module(), "NOPKernel"); + hipFunction_t f = GetKernel(GetModule(), "NOPKernel"); HIP_CHECK(func(f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, 0, nullptr, nullptr, nullptr, nullptr, nullptr, 0u)); HIP_CHECK(hipDeviceSynchronize()); @@ -120,7 +119,7 @@ template void ModuleLaunchKernelPositiveParamet template void ModuleLaunchKernelNegativeParameters( bool extLaunch = false) { - hipFunction_t f = GetKernel(mg.module(), "NOPKernel"); + hipFunction_t f = GetKernel(GetModule(), "NOPKernel"); hipError_t expectedErrorLaunchParam = (extLaunch == true) ? hipErrorInvalidConfiguration : hipErrorInvalidValue; hipError_t expectedErrorOverCapacityGridDim = (extLaunch == true) ? hipSuccess @@ -213,7 +212,7 @@ template void ModuleLaunchKernelNegativeParamet } SECTION("Passing kernel_args and extra simultaneously") { - hipFunction_t f = GetKernel(mg.module(), "Kernel42"); + hipFunction_t f = GetKernel(GetModule(), "Kernel42"); LinearAllocGuard result_dev(LinearAllocs::hipMalloc, sizeof(int)); int* result_ptr = result_dev.ptr(); size_t size = sizeof(result_ptr); @@ -230,7 +229,7 @@ template void ModuleLaunchKernelNegativeParamet } SECTION("Invalid extra") { - hipFunction_t f = GetKernel(mg.module(), "Kernel42"); + hipFunction_t f = GetKernel(GetModule(), "Kernel42"); void* extra[0] = {}; HIP_CHECK_ERROR(func(f, 1, 1, 1, 1, 1, 1, 0, nullptr, nullptr, extra, nullptr, nullptr, 0u), hipErrorInvalidValue); diff --git a/catch/unit/module/not_a_module.txt b/catch/unit/module/not_a_module.txt index e69de29bb..5a97d7c38 100644 --- a/catch/unit/module/not_a_module.txt +++ b/catch/unit/module/not_a_module.txt @@ -0,0 +1 @@ +This is not a module!