From c1764bd1f1644a146715ca97b3adc761255ef53d Mon Sep 17 00:00:00 2001 From: Paulius Velesko Date: Sat, 9 Mar 2024 07:02:16 +0200 Subject: [PATCH 01/18] Integrate chipstar --- catch/CMakeLists.txt | 25 +++++++++++++------ .../external/Catch2/cmake/Catch2/Catch.cmake | 18 ++++++++++++- catch/hipTestMain/CMakeLists.txt | 5 +++- catch/include/hip_test_context.hh | 6 +++++ catch/kernels/CMakeLists.txt | 5 ++++ 5 files changed, 50 insertions(+), 9 deletions(-) diff --git a/catch/CMakeLists.txt b/catch/CMakeLists.txt index 7766df816..bbbd15b87 100644 --- a/catch/CMakeLists.txt +++ b/catch/CMakeLists.txt @@ -12,7 +12,7 @@ option(ENABLE_ADDRESS_SANITIZER "Option to enable ASAN build" OFF) message(STATUS "STANDALONE_TESTS : ${STANDALONE_TESTS}") # Check if platform is set -if(NOT HIP_PLATFORM STREQUAL "amd" AND NOT HIP_PLATFORM STREQUAL "nvidia") +if(NOT HIP_PLATFORM STREQUAL "amd" AND NOT HIP_PLATFORM STREQUAL "nvidia" AND NOT HIP_PLATFORM STREQUAL "spirv") message(FATAL_ERROR "Unexpected HIP_PLATFORM: " ${HIP_PLATFORM}) endif() @@ -192,7 +192,17 @@ message(STATUS "CMAKE HIP ARCHITECTURES: ${CMAKE_HIP_ARCHITECTURES}") # preference to pass arch - # OFFLOAD_ARCH_STR # rocm_agent_enumerator -if(NOT DEFINED OFFLOAD_ARCH_STR +if(HIP_PLATFORM STREQUAL "spirv" AND NOT MASTER_PROJECT) + # for CHIP-SPV, OFFLOAD_ARCH_STR will be passed to CMAKE_CXX_FLAGS which will make + # compile-only flags (such as -x hip) to be passed to linker. This will cause strange errors. + # These could be removed, but then we will need to add manual linking of libCHIP.so so + # might as well just use a different approach for CHIP-SPV path. + # set(OFFLOAD_ARCH_STR ${HIP_OFFLOAD_COMPILE_OPTIONS_BUILD}) + message(WARNING "HIP_PLATFORM is spirv, OFFLOAD_ARCH_STR is set to ${OFFLOAD_ARCH_STR}") +elseif(HIP_PLATFORM STREQUAL "spirv" AND MASTER_PROJECT) + # TODO skip setting this & enforce use of hipcc + message(FATAL_ERROR "building hip-tests with HIP_PLATFORM=spirv is not supported when building as a standalone project") +elseif(NOT DEFINED OFFLOAD_ARCH_STR AND EXISTS "${ROCM_PATH}/bin/rocm_agent_enumerator" AND HIP_PLATFORM STREQUAL "amd" AND UNIX) execute_process(COMMAND "${ROCM_PATH}/bin/rocm_agent_enumerator" @@ -297,11 +307,12 @@ add_subdirectory(unit ${CATCH_BUILD_DIR}/unit) add_subdirectory(ABM ${CATCH_BUILD_DIR}/ABM) add_subdirectory(kernels ${CATCH_BUILD_DIR}/kernels) add_subdirectory(hipTestMain ${CATCH_BUILD_DIR}/hipTestMain) -add_subdirectory(stress ${CATCH_BUILD_DIR}/stress) -add_subdirectory(TypeQualifiers ${CATCH_BUILD_DIR}/TypeQualifiers) -add_subdirectory(perftests ${CATCH_BUILD_DIR}/perftests) -add_subdirectory(multiproc ${CATCH_BUILD_DIR}/multiproc) -add_subdirectory(performance ${CATCH_BUILD_DIR}/performance) +# add_subdirectory(stress ${CATCH_BUILD_DIR}/stress) +# add_subdirectory(TypeQualifiers ${CATCH_BUILD_DIR}/TypeQualifiers) +# add_subdirectory(perftests ${CATCH_BUILD_DIR}/perftests) +# add_subdirectory(multiproc ${CATCH_BUILD_DIR}/multiproc) +# add_subdirectory(performance ${CATCH_BUILD_DIR}/performance) + cmake_policy(POP) diff --git a/catch/external/Catch2/cmake/Catch2/Catch.cmake b/catch/external/Catch2/cmake/Catch2/Catch.cmake index c6c747155..e795b5635 100644 --- a/catch/external/Catch2/cmake/Catch2/Catch.cmake +++ b/catch/external/Catch2/cmake/Catch2/Catch.cmake @@ -358,6 +358,8 @@ function(hip_add_exe_to_target) get_filename_component(_EXE_NAME ${SRC_NAME} NAME_WLE) endif() + + # Create shared lib of all tests if(NOT RTC_TESTING) add_executable(${_EXE_NAME} EXCLUDE_FROM_ALL ${SRC_NAME} ${COMMON_SHARED_SRC} $ $) @@ -365,8 +367,13 @@ function(hip_add_exe_to_target) add_executable(${_EXE_NAME} EXCLUDE_FROM_ALL ${SRC_NAME} ${COMMON_SHARED_SRC} $) if(HIP_PLATFORM STREQUAL "amd") target_link_libraries(${_EXE_NAME} hiprtc) - else() + elseif(HIP_PLATFORM STREQUAL "nvidia") target_link_libraries(${_EXE_NAME} nvrtc) + elseif(HIP_PLATFORM STREQUAL "spirv") + message(FATAL_ERROR "RTC path for SPIRV not yet checked") + target_link_libraries(${_EXE_NAME} spirv) + else() + message(FATAL_ERROR "Unsupported HIP_PLATFORM: ${HIP_PLATFORM}") endif() endif() if (DEFINED _PROPERTY) @@ -390,6 +397,15 @@ function(hip_add_exe_to_target) target_link_libraries(${_EXE_NAME} ${_LINKER_LIBS}) endif() + # link against CHIP-SPV + # Required because CHIP-SPV path is not using OFFLOAD_ARCH_STR + # see hip-tests/CMakeLists.txt + if(NOT MASTER_PROJECT AND HIP_PLATFORM STREQUAL "spirv") + # target_compile_options(${_EXE_NAME} PRIVATE -mllvm -amdgpu-early-inline-all=true) + target_compile_options(${_EXE_NAME} PRIVATE ${HIP_OFFLOAD_COMPILE_OPTIONS_BUILD_}) + target_link_libraries(${_EXE_NAME} CHIP) + endif() + # Add dependency on build_tests to build it on this custom target add_dependencies(${_TEST_TARGET_NAME} ${_EXE_NAME}) diff --git a/catch/hipTestMain/CMakeLists.txt b/catch/hipTestMain/CMakeLists.txt index 95b7a0954..1be9cf082 100644 --- a/catch/hipTestMain/CMakeLists.txt +++ b/catch/hipTestMain/CMakeLists.txt @@ -25,6 +25,9 @@ endif() add_library(Main_Object EXCLUDE_FROM_ALL OBJECT main.cc hip_test_context.cc hip_test_features.cc) if(HIP_PLATFORM MATCHES "amd") set_property(TARGET Main_Object PROPERTY CXX_STANDARD 17) -else() +elseif(HIP_PLATFORM MATCHES "nvidia") target_compile_options(Main_Object PUBLIC -std=c++17) +elseif(HIP_PLATFORM MATCHES "spirv") + target_compile_options(Main_Object PUBLIC ${HIP_OFFLOAD_COMPILE_OPTIONS_BUILD_}) + set_property(TARGET Main_Object PROPERTY CXX_STANDARD 17) endif() diff --git a/catch/include/hip_test_context.hh b/catch/include/hip_test_context.hh index 8e06c3fbb..af603c8b1 100644 --- a/catch/include/hip_test_context.hh +++ b/catch/include/hip_test_context.hh @@ -47,9 +47,15 @@ THE SOFTWARE. #if defined(__HIP_PLATFORM_AMD__) #define HT_AMD 1 #define HT_NVIDIA 0 +#define HT_SPIRV 0 #elif defined(__HIP_PLATFORM_NVIDIA__) #define HT_AMD 0 #define HT_NVIDIA 1 +#define HT_SPIRV 0 +#elif defined(__HIP_PLATFORM_CLANG__) || defined(__HIP_PLATFORM_SPIRV__) +#define HT_AMD 0 +#define HT_NVIDIA 0 +#define HT_SPIRV 1 #else #error "Platform not recognized" #endif diff --git a/catch/kernels/CMakeLists.txt b/catch/kernels/CMakeLists.txt index 91e1ab69f..8b346dcb2 100644 --- a/catch/kernels/CMakeLists.txt +++ b/catch/kernels/CMakeLists.txt @@ -5,4 +5,9 @@ if(NOT RTC_TESTING) add_library(KERNELS EXCLUDE_FROM_ALL OBJECT ${TEST_SRC}) target_compile_options(KERNELS PUBLIC -std=c++17) + + # If compiling as part of CHIP-SPV, add the necessary offload flags which would normally be added by hipcc + if(HIP_PLATFORM STREQUAL "spirv" AND NOT MASTER_PROJECT) + target_compile_options(KERNELS PUBLIC ${HIP_OFFLOAD_COMPILE_OPTIONS_BUILD_}) + endif() endif() From 3787a78a00c71f48cec566e20c47c18a5fb3c8c3 Mon Sep 17 00:00:00 2001 From: Paulius Velesko Date: Mon, 9 Jan 2023 08:20:02 +0000 Subject: [PATCH 02/18] CHIP-SPV Integration contexts --- catch/hipTestMain/hip_test_context.cc | 11 ++++++++++- catch/include/hip_test_common.hh | 2 +- catch/include/hip_test_context.hh | 5 +++-- 3 files changed, 14 insertions(+), 4 deletions(-) diff --git a/catch/hipTestMain/hip_test_context.cc b/catch/hipTestMain/hip_test_context.cc index 390da2545..f80c8ce1e 100644 --- a/catch/hipTestMain/hip_test_context.cc +++ b/catch/hipTestMain/hip_test_context.cc @@ -19,6 +19,8 @@ void TestContext::detectOS() { void TestContext::detectPlatform() { #if (HT_AMD == 1) amd = true; +#elif (HT_SPIRV == 1) + spirv = true; #elif (HT_NVIDIA == 1) nvidia = true; #endif @@ -160,7 +162,13 @@ std::string& TestContext::getCommonJsonFile() { void TestContext::getConfigFiles() { - config_.platform = (amd ? "amd" : (nvidia ? "nvidia" : "unknown")); + if(config_.platform == "amd") { + amd = true; + } else if(config_.platform == "nvidia") { + nvidia = true; + } else if(config_.platform == "spirv") { + spirv = true; + } config_.os = (p_windows ? "windows" : (p_linux ? "linux" : "unknown")); if (config_.os == "unknown" || config_.platform == "unknown") { @@ -210,6 +218,7 @@ bool TestContext::isLinux() const { return p_linux; } bool TestContext::isNvidia() const { return nvidia; } bool TestContext::isAmd() const { return amd; } +bool TestContext::isSpirv() const { return spirv; } void TestContext::parseOptions(int argc, char** argv) { // Test name is at [1] position diff --git a/catch/include/hip_test_common.hh b/catch/include/hip_test_common.hh index 21707f761..7786147e4 100644 --- a/catch/include/hip_test_common.hh +++ b/catch/include/hip_test_common.hh @@ -254,7 +254,7 @@ static inline int RAND_R(unsigned* rand_seed) { inline bool isImageSupported() { int imageSupport = 1; -#if HT_AMD +#if HT_AMD || HT_SPIRV int device; HIP_CHECK(hipGetDevice(&device)); HIPCHECK(hipDeviceGetAttribute(&imageSupport, hipDeviceAttributeImageSupport, device)); diff --git a/catch/include/hip_test_context.hh b/catch/include/hip_test_context.hh index af603c8b1..9dca6a49c 100644 --- a/catch/include/hip_test_context.hh +++ b/catch/include/hip_test_context.hh @@ -80,12 +80,12 @@ struct HCResult { class TestContext { bool p_windows = false, p_linux = false; // OS - bool amd = false, nvidia = false; // HIP Platform + bool amd = false, nvidia = false, spirv = false; // HIP Platform std::string exe_path; std::string current_test; std::set skip_test; std::string json_file_; - std::vector platform_list_ = {"amd", "nvidia"}; + std::vector platform_list_ = {"amd", "nvidia", "spirv"}; std::vector os_list_ = {"windows", "linux", "all"}; std::vector amd_arch_list_ = {}; @@ -147,6 +147,7 @@ class TestContext { bool isLinux() const; bool isNvidia() const; bool isAmd() const; + bool isSpirv() const; bool skipTest() const; const std::string& getCurrentTest() const { return current_test; } From 7669fa879f84bcaacb1e960cce84dbc174366054 Mon Sep 17 00:00:00 2001 From: Paulius Velesko Date: Sat, 9 Mar 2024 08:38:03 +0200 Subject: [PATCH 03/18] HT_AMD == HT_SPIRV where the same --- catch/include/memcpy3d_tests_common.hh | 8 +++--- catch/include/utils.hh | 4 +-- catch/unit/assertion/assert.cc | 4 +-- .../errorHandling/error_handling_common.cc | 28 +++++++++---------- 4 files changed, 22 insertions(+), 22 deletions(-) diff --git a/catch/include/memcpy3d_tests_common.hh b/catch/include/memcpy3d_tests_common.hh index e55469534..fc9f06e18 100644 --- a/catch/include/memcpy3d_tests_common.hh +++ b/catch/include/memcpy3d_tests_common.hh @@ -595,7 +595,7 @@ void Memcpy3DZeroWidthHeightDepth(F memcpy_func, const hipStream_t stream = null } constexpr auto MemTypeHost() { -#if HT_AMD +#if HT_AMD || HT_SPIRV return hipMemoryTypeHost; #else return CU_MEMORYTYPE_HOST; @@ -603,7 +603,7 @@ constexpr auto MemTypeHost() { } constexpr auto MemTypeDevice() { -#if HT_AMD +#if HT_AMD || HT_SPIRV return hipMemoryTypeDevice; #else return CU_MEMORYTYPE_DEVICE; @@ -611,7 +611,7 @@ constexpr auto MemTypeDevice() { } constexpr auto MemTypeArray() { -#if HT_AMD +#if HT_AMD || HT_SPIRV return hipMemoryTypeArray; #else return CU_MEMORYTYPE_ARRAY; @@ -619,7 +619,7 @@ constexpr auto MemTypeArray() { } constexpr auto MemTypeUnified() { -#if HT_AMD +#if HT_AMD || HT_SPIRV return hipMemoryTypeUnified; #else return CU_MEMORYTYPE_UNIFIED; diff --git a/catch/include/utils.hh b/catch/include/utils.hh index 3855308a4..457d215bf 100644 --- a/catch/include/utils.hh +++ b/catch/include/utils.hh @@ -128,7 +128,7 @@ static __global__ void Delay(uint32_t interval, const uint32_t ticks_per_ms) { __builtin_amdgcn_s_sleep(10); } #endif - #if HT_NVIDIA + #if HT_NVIDIA || HT_SPIRV uint64_t start = clock64(); while (clock64() - start < ticks_per_ms) { } @@ -150,7 +150,7 @@ __global__ void Iota(T* const out, size_t pitch, size_t w, size_t h, size_t d) { inline void LaunchDelayKernel(const std::chrono::milliseconds interval, const hipStream_t stream = nullptr) { int ticks_per_ms = 0; - #if HT_AMD + #if HT_AMD || HT_SPIRV HIPCHECK(hipDeviceGetAttribute(&ticks_per_ms, hipDeviceAttributeWallClockRate, 0)); #endif #if HT_NVIDIA diff --git a/catch/unit/assertion/assert.cc b/catch/unit/assertion/assert.cc index 29cadd289..b27fb4976 100644 --- a/catch/unit/assertion/assert.cc +++ b/catch/unit/assertion/assert.cc @@ -68,7 +68,7 @@ template void LaunchAssertKernel() { if constexpr (should_abort) { AssertFailKernel<<>>(d_a); -#if HT_AMD +#if HT_AMD || HT_SPIRV HIP_CHECK(hipDeviceSynchronize()); #else HIP_CHECK_ERROR(hipDeviceSynchronize(), hipErrorAssert); @@ -116,7 +116,7 @@ TEST_CASE("Unit_Assert_Positive_Basic_KernelPass") { */ TEST_CASE("Unit_Assert_Positive_Basic_KernelFail") { try_and_catch_abort(&LaunchAssertKernel); -#if HT_AMD +#if HT_AMD || HT_SPIRV REQUIRE(abort_raised_flag == 1); #else REQUIRE(abort_raised_flag == 0); diff --git a/catch/unit/errorHandling/error_handling_common.cc b/catch/unit/errorHandling/error_handling_common.cc index 20267e793..879538002 100644 --- a/catch/unit/errorHandling/error_handling_common.cc +++ b/catch/unit/errorHandling/error_handling_common.cc @@ -24,7 +24,7 @@ THE SOFTWARE. const char* ErrorName(hipError_t enumerator) { switch (enumerator) { -#if HT_AMD +#if HT_AMD || HT_SPIRV case hipSuccess: return "hipSuccess"; case hipErrorInvalidValue: @@ -343,7 +343,7 @@ const char* ErrorString(hipError_t enumerator) { case hipErrorProfilerDisabled: return "profiler disabled while using external profiling tool"; case hipErrorProfilerNotInitialized: -#if HT_AMD +#if HT_AMD || HT_SPIRV return "profiler is not initialized"; #elif HT_NVIDIA return "profiler not initialized: call cudaProfilerInitialize()"; @@ -352,62 +352,62 @@ const char* ErrorString(hipError_t enumerator) { return "profiler already started"; case hipErrorProfilerAlreadyStopped: return "profiler already stopped"; -#if HT_AMD +#if HT_AMD || HT_SPIRV case hipErrorInvalidConfiguration: return "invalid configuration argument"; #elif HT_NVIDIA return "unknown error"; #endif -#if HT_AMD +#if HT_AMD || HT_SPIRV case hipErrorInvalidPitchValue: return "invalid pitch argument"; #elif HT_NVIDIA return "unknown error"; #endif -#if HT_AMD +#if HT_AMD || HT_SPIRV case hipErrorInvalidSymbol: return "invalid device symbol"; #elif HT_NVIDIA return "unknown error"; #endif -#if HT_AMD +#if HT_AMD || HT_SPIRV case hipErrorInvalidDevicePointer: return "invalid device pointer"; #elif HT_NVIDIA return "unknown error"; #endif -#if HT_AMD +#if HT_AMD || HT_SPIRV case hipErrorInvalidMemcpyDirection: return "invalid copy direction for memcpy"; #elif HT_NVIDIA return "unknown error"; #endif -#if HT_AMD +#if HT_AMD || HT_SPIRV case hipErrorInsufficientDriver: return "driver version is insufficient for runtime version"; #elif HT_NVIDIA return "unknown error"; #endif -#if HT_AMD +#if HT_AMD || HT_SPIRV case hipErrorMissingConfiguration: return "__global__ function call is not configured"; #elif HT_NVIDIA return "unknown error"; #endif -#if HT_AMD +#if HT_AMD || HT_SPIRV case hipErrorPriorLaunchFailure: return "unspecified launch failure in prior launch"; #elif HT_NVIDIA return "unknown error"; #endif -#if HT_AMD +#if HT_AMD || HT_SPIRV case hipErrorInvalidDeviceFunction: return "invalid device function"; #elif HT_NVIDIA return "unknown error"; #endif case hipErrorNoDevice: -#if HT_AMD +#if HT_AMD || HT_SPIRV return "no ROCm-capable device is detected"; #elif HT_NVIDIA return "no CUDA-capable device is detected"; @@ -419,7 +419,7 @@ const char* ErrorString(hipError_t enumerator) { case hipErrorInvalidContext: return "invalid device context"; case hipErrorContextAlreadyCurrent: -#if HT_AMD +#if HT_AMD || HT_SPIRV return "context is already current context"; #elif HT_NVIDIA return "context already current"; @@ -451,7 +451,7 @@ const char* ErrorString(hipError_t enumerator) { case hipErrorPeerAccessUnsupported: return "peer access is not supported between these two devices"; case hipErrorInvalidKernelFile: -#if HT_AMD +#if HT_AMD || HT_SPIRV return "invalid kernel file"; #elif HT_NVIDIA return "a PTX JIT compilation failed"; From 9ac99d721cf65cae994cae5262950ee3bfde5ba0 Mon Sep 17 00:00:00 2001 From: Paulius Velesko Date: Sat, 9 Mar 2024 08:38:32 +0200 Subject: [PATCH 04/18] Oversubscription tests skip on SPIRV --- catch/stress/memory/hipHmmOvrSubscriptionTst.cc | 3 +++ catch/stress/memory/hipHostRegisterStress.cc | 3 +++ 2 files changed, 6 insertions(+) diff --git a/catch/stress/memory/hipHmmOvrSubscriptionTst.cc b/catch/stress/memory/hipHmmOvrSubscriptionTst.cc index 6c7abf210..1bab3e5b2 100644 --- a/catch/stress/memory/hipHmmOvrSubscriptionTst.cc +++ b/catch/stress/memory/hipHmmOvrSubscriptionTst.cc @@ -35,6 +35,9 @@ __global__ void floatx2(float* ptr, size_t size) { } TEST_CASE("Stress_HMM_OverSubscriptionTst") { +#if HT_SPIRV + HipTest::HIP_SKIP_TEST("Stress_HMM_OverSubscriptionTst Unsupported on SPIRV"); +#endif int hmm = 0; HIP_CHECK(hipDeviceGetAttribute(&hmm, hipDeviceAttributeManagedMemory, 0)); diff --git a/catch/stress/memory/hipHostRegisterStress.cc b/catch/stress/memory/hipHostRegisterStress.cc index 564dfc29b..dc53a8306 100644 --- a/catch/stress/memory/hipHostRegisterStress.cc +++ b/catch/stress/memory/hipHostRegisterStress.cc @@ -56,6 +56,9 @@ static __global__ void Inc(uint8_t* Ad) { * - HIP_VERSION >= 5.6 */ TEST_CASE("Stress_hipHostRegister_Oversubscription") { +#if HT_SPIRV + HipTest::HIP_SKIP_TEST("Stress_hipHostRegister_Oversubscription Unsupported on SPIRV"); +#endif hipDeviceProp_t prop; HIP_CHECK(hipGetDeviceProperties(&prop, 0)); std::string arch = prop.gcnArchName; From 9f199ca6e9a4fe43734255502b4a8756d37289df Mon Sep 17 00:00:00 2001 From: Paulius Velesko Date: Sun, 10 Mar 2024 08:06:38 +0200 Subject: [PATCH 05/18] Add fix for valideArguments --- catch/include/hip_test_common.hh | 53 ++++++++++++++++++++++++++------ 1 file changed, 43 insertions(+), 10 deletions(-) diff --git a/catch/include/hip_test_common.hh b/catch/include/hip_test_common.hh index 7786147e4..96e7b27e4 100644 --- a/catch/include/hip_test_common.hh +++ b/catch/include/hip_test_common.hh @@ -276,7 +276,9 @@ static inline void HIP_SKIP_TEST(char const* const reason) noexcept { * * @return constexpr std::tuple the expected arguments of the kernel. */ -template std::tuple getExpectedArgs(void(FArgs...)){}; +// template std::tuple getExpectedArgs(void(FArgs...)){}; +template +std::tuple getExpectedArgs(void(*)(FArgs...)) {}; /** * @brief Asserts that the types of the arguments of a function match exactly with the types in the @@ -289,10 +291,18 @@ template std::tuple getExpectedArgs(void(FArgs...) * @tparam F the kernel function * @tparam Args the parameters that will be passed to the kernel. */ -template void validateArguments(F f, Args...) { - using expectedArgsTuple = decltype(getExpectedArgs(f)); - static_assert(std::is_same>::value, - "Kernel arguments types must match exactly!"); +// template void validateArguments(F f, Args...) { +// using expectedArgsTuple = decltype(getExpectedArgs(f)); +// static_assert(std::is_same>::value, +// "Kernel arguments types must match exactly!"); +// } +template +void validateArguments(F f, Args&&... args) { + using expectedArgsTuple = decltype(getExpectedArgs(f)); + using providedArgsTuple = std::tuple; + + static_assert(std::is_same::value, + "Kernel arguments types must match exactly!"); } /** @@ -311,12 +321,35 @@ template void validateArguments(F f, Args...) { * @param stream * @param packedArgs A list of kernel arguments to be forwarded. */ -template -void launchKernel(K kernel, Dim numBlocks, Dim numThreads, std::uint32_t memPerBlock, - hipStream_t stream, Args&&... packedArgs) { +// template +// void launchKernel(K kernel, Dim numBlocks, Dim numThreads, std::uint32_t memPerBlock, +// hipStream_t stream, Args&&... packedArgs) { +// #ifndef RTC_TESTING +// validateArguments(kernel, packedArgs...); +// kernel<<>>(std::forward(packedArgs)...); +// #else +// launchRTCKernel(kernel, numBlocks, numThreads, memPerBlock, stream, +// std::forward(packedArgs)...); +// #endif +// HIP_CHECK(hipGetLastError()); +// } + +template +void launchKernel(Kernel kernel, Dim numBlocks, Dim numThreads, std::uint32_t memPerBlock, hipStream_t stream, Args&&... args) { #ifndef RTC_TESTING - validateArguments(kernel, packedArgs...); - kernel<<>>(std::forward(packedArgs)...); + // Define a stateless, capture-free lambda that matches the kernel's signature. + auto kernelWrapperLambda = [] (Args... args) { + // This lambda is intentionally left empty as it's used solely for type validation. + }; + + // Convert the lambda to a function pointer. + void (*kernelWrapper)(Args...) = kernelWrapperLambda; + + // Use the wrapper function pointer to validate arguments. + validateArguments(kernelWrapper, std::forward(args)...); + + // Launch the kernel directly with the provided arguments. + kernel<<>>(std::forward(args)...); #else launchRTCKernel(kernel, numBlocks, numThreads, memPerBlock, stream, std::forward(packedArgs)...); From 7678b4374f56853df7d3c910d2a32ca9ab656925 Mon Sep 17 00:00:00 2001 From: Paulius Velesko Date: Mon, 11 Mar 2024 02:55:41 +0200 Subject: [PATCH 06/18] Modifications to copile on chipStar --- catch/packaging/CMakeLists.txt | 1 + catch/unit/CMakeLists.txt | 14 +++++------ catch/unit/deviceLib/CMakeLists.txt | 23 +++++++++++-------- catch/unit/event/Unit_hipEventRecord.cc | 4 ++-- .../graph_memcpy_to_from_symbol_common.hh | 2 +- catch/unit/kernel/CMakeLists.txt | 2 +- catch/unit/memory/CMakeLists.txt | 9 +++++--- catch/unit/memory/hipArray3DCreate.cc | 8 +++---- catch/unit/memory/hipArrayCommon.hh | 2 +- catch/unit/memory/hipArrayCreate.cc | 2 +- catch/unit/memory/hipGetSymbolSizeAddress.cc | 6 ++--- catch/unit/memory/hipMalloc3DArray.cc | 12 +++++----- catch/unit/memory/hipMallocArray.cc | 16 ++++++------- catch/unit/memory/hipMallocMipmappedArray.cc | 10 ++++---- catch/unit/memory/hipMemcpyAsync.cc | 8 +++---- catch/unit/memory/hipPointerGetAttribute.cc | 8 +++---- catch/unit/module/CMakeLists.txt | 4 ++-- catch/unit/printf/CMakeLists.txt | 15 ++++++------ catch/unit/stream/hipStreamACb_MultiThread.cc | 2 +- .../stream/hipStreamACb_StrmSyncTiming.cc | 2 +- catch/unit/stream/hipStreamAddCallback.cc | 2 +- catch/unit/streamperthread/CMakeLists.txt | 4 ++-- catch/unit/texture/CMakeLists.txt | 4 ++-- catch/unit/warp/CMakeLists.txt | 6 ++--- 24 files changed, 87 insertions(+), 79 deletions(-) diff --git a/catch/packaging/CMakeLists.txt b/catch/packaging/CMakeLists.txt index 1b124bb84..5a23982fb 100644 --- a/catch/packaging/CMakeLists.txt +++ b/catch/packaging/CMakeLists.txt @@ -112,4 +112,5 @@ set(CPACK_TEST_ZIP "ON") set(CPACK_ZIP_TEST_PACKAGE_NAME "catch") endif() +set(CPACK_RESOURCE_FILE_LICENSE "${CMAKE_CURRENT_SOURCE_DIR}/../../LICENSE.txt") include(CPack) diff --git a/catch/unit/CMakeLists.txt b/catch/unit/CMakeLists.txt index 6b63292c9..4fa382c63 100644 --- a/catch/unit/CMakeLists.txt +++ b/catch/unit/CMakeLists.txt @@ -25,25 +25,25 @@ add_subdirectory(memory) add_subdirectory(stream_ordered) add_subdirectory(stream) add_subdirectory(event) -add_subdirectory(occupancy) +# add_subdirectory(occupancy) tex2gather add_subdirectory(device) add_subdirectory(printf) -add_subdirectory(texture) -add_subdirectory(surface) +# add_subdirectory(texture) +# add_subdirectory(surface) add_subdirectory(streamperthread) add_subdirectory(kernel) add_subdirectory(multiThread) add_subdirectory(compiler) add_subdirectory(errorHandling) -add_subdirectory(cooperativeGrps) +# add_subdirectory(cooperativeGrps) add_subdirectory(context) add_subdirectory(device_memory) add_subdirectory(warp) -add_subdirectory(dynamicLoading) +# add_subdirectory(dynamicLoading) # not supported by chipstar? add_subdirectory(g++) -add_subdirectory(module) +# add_subdirectory(module) add_subdirectory(channelDescriptor) -add_subdirectory(executionControl) +# add_subdirectory(executionControl) # cooperative_groups add_subdirectory(math) add_subdirectory(vector_types) add_subdirectory(atomics) diff --git a/catch/unit/deviceLib/CMakeLists.txt b/catch/unit/deviceLib/CMakeLists.txt index f93a7a43e..1d7b746f0 100644 --- a/catch/unit/deviceLib/CMakeLists.txt +++ b/catch/unit/deviceLib/CMakeLists.txt @@ -33,7 +33,6 @@ set(TEST_SRC syncthreadsand.cc syncthreadscount.cc syncthreadsor.cc - deviceAllocation.cc Atomic_func.cc DoublePrecisionIntrinsics.cc DoublePrecisionMathDevice.cc @@ -46,26 +45,28 @@ set(TEST_SRC SimpleAtomicsTest.cc hipTestAtomicAdd.cc hipStdComplex.cc - hipTestClock.cc hip_trig.cc hipDeviceMemcpy.cc hipTestIncludeMath.cc hipTestDotFunctions.cc hipTestDeviceSymbol.cc - hipTestNew.cc hipThreadFence.cc hipTestDevice.cc hipTestDeviceLimit.cc - hipTestDeviceDouble.cc - hipTestHost.cc ) if(HIP_PLATFORM MATCHES "nvidia") set_source_files_properties(hipTestHost.cc PROPERTIES COMPILE_OPTIONS "--expt-relaxed-constexpr") endif() -if(UNIX) +if(UNIX AND NOT HIP_PLATFORM MATCHES "spirv") set(TEST_SRC ${TEST_SRC} - deviceAllocation.cc) + deviceAllocation.cc # unsupported device-side malloc + hipTestNew.cc # unsupported device-side new + hipTestClock.cc # unsupported __clock() and __clock64() + hipTestDeviceDouble.cc # SPIR-V Translator: InvalidBitWidth: Invalid bit width in input: 128 + hipTestHost.cc # SPIR-V Translator: InvalidBitWidth: Invalid bit width in input: 128 + ) + endif() # AMD only tests @@ -181,11 +182,13 @@ endif() TEST_SRC ${TEST_SRC} TEST_TARGET_NAME build_tests LINKER_LIBS hiprtc) -elseif(HIP_PLATFORM MATCHES "nvidia") +elseif(HIP_PLATFORM MATCHES "nvidia" OR HIP_PLATFORM MATCHES "spirv") hip_add_exe_to_target(NAME UnitDeviceTests TEST_SRC ${TEST_SRC} TEST_TARGET_NAME build_tests - COMPILE_OPTIONS --Wno-deprecated-declarations) + COMPILE_OPTIONS -Wno-deprecated-declarations) #--Wno-deprecated-declarations unrecognized clang++ endif() -add_dependencies(build_tests kerDevAllocMultCO.code kerDevWriteMultCO.code kerDevFreeMultCO.code kerDevAllocSingleKer.code) +if (NOT HIP_PLATFORM MATCHES "spirv") + add_dependencies(build_tests kerDevAllocMultCO.code kerDevWriteMultCO.code kerDevFreeMultCO.code kerDevAllocSingleKer.code) +endif() diff --git a/catch/unit/event/Unit_hipEventRecord.cc b/catch/unit/event/Unit_hipEventRecord.cc index 7dd8c582b..e20f39908 100644 --- a/catch/unit/event/Unit_hipEventRecord.cc +++ b/catch/unit/event/Unit_hipEventRecord.cc @@ -69,14 +69,14 @@ TEST_CASE("Unit_hipEventRecord") { WithFlags_Default = hipEventDefault, WithFlags_Blocking = hipEventBlockingSync, WithFlags_DisableTiming = hipEventDisableTiming, -#if HT_AMD +#if HT_AMD || HT_SPIRV WithFlags_ReleaseToDevice = hipEventReleaseToDevice, WithFlags_ReleaseToSystem = hipEventReleaseToSystem, #endif WithoutFlags }; -#if HT_AMD +#if HT_AMD || HT_SPIRV auto flags = GENERATE(WithFlags_Default, WithFlags_Blocking, WithFlags_DisableTiming, WithFlags_ReleaseToDevice, WithFlags_ReleaseToSystem, WithoutFlags); #endif diff --git a/catch/unit/graph/graph_memcpy_to_from_symbol_common.hh b/catch/unit/graph/graph_memcpy_to_from_symbol_common.hh index e2c183b78..689d3adf9 100644 --- a/catch/unit/graph/graph_memcpy_to_from_symbol_common.hh +++ b/catch/unit/graph/graph_memcpy_to_from_symbol_common.hh @@ -163,7 +163,7 @@ void MemcpyToSymbolCommonNegative(F f, const void* symbol, void* src, size_t cou #endif } -#if HT_AMD +#if HT_AMD || HT_SPIRV #define SYMBOL(expr) &HIP_SYMBOL(expr) #else #define SYMBOL(expr) HIP_SYMBOL(expr) diff --git a/catch/unit/kernel/CMakeLists.txt b/catch/unit/kernel/CMakeLists.txt index 2e7c0aecf..ec6f3997a 100644 --- a/catch/unit/kernel/CMakeLists.txt +++ b/catch/unit/kernel/CMakeLists.txt @@ -22,7 +22,7 @@ set(TEST_SRC hipMemFaultStackAllocation.cc hipLaunchBounds.cc - hipShflTests.cc + # hipShflTests.cc hipDynamicShared.cc hipDynamicShared2.cc hipEmptyKernel.cc diff --git a/catch/unit/memory/CMakeLists.txt b/catch/unit/memory/CMakeLists.txt index 025e369b5..28cb545f7 100644 --- a/catch/unit/memory/CMakeLists.txt +++ b/catch/unit/memory/CMakeLists.txt @@ -101,7 +101,7 @@ if(HIP_PLATFORM MATCHES "amd") set(TEST_SRC ${TEST_SRC} hipPointerSetAttribute.cc) endif() else() - set(TEST_SRC ${TEST_SRC} hipGetSymbolSizeAddress.cc) + # set(TEST_SRC ${TEST_SRC} hipGetSymbolSizeAddress.cc) endif() hip_add_exe_to_target(NAME MemoryTest1 @@ -136,6 +136,8 @@ set(TEST_SRC hipMemcpyDtoD.cc hipMemcpyDtoDAsync.cc hipHostMalloc.cc + # /space/pvelesko/chipStar/hip-tests-update/hip-tests/catch/./include/hip_test_common.hh:296:38: error: no matching function for call to 'getExpectedArgs' + # 296 | using expectedArgsTuple = decltype(getExpectedArgs(f)); hipMemcpy_old.cc hipMemcpy_derivatives.cc hipMemcpyAsync.cc @@ -164,7 +166,7 @@ set(TEST_SRC hipMemRangeGetAttributes.cc hipFreeAsync.cc hipMallocAsync.cc - hipStreamAttachMemAsync.cc + # hipStreamAttachMemAsync.cc # getExpectedArgs hipMemRangeGetAttributes_old.cc hipMemGetAddressRange.cc hipMallocMipmappedArray.cc @@ -179,7 +181,8 @@ if(HIP_PLATFORM MATCHES "amd") hipArray3DGetDescriptor.cc) endif() -set(NOT_FOR_MI200_AND_ABOVE_TEST hipMallocArray.cc hipArrayCreate.cc) # tests not for MI200+ +# set(NOT_FOR_MI200_AND_ABOVE_TEST hipMallocArray.cc hipArrayCreate.cc) # tests not for MI200+ +set(NOT_FOR_MI200_AND_ABOVE_TEST ) # tests not for MI200+ set(MI200_AND_ABOVE_TARGETS gfx90a gfx940 gfx941 gfx942) function(CheckRejectedArchs OFFLOAD_ARCH_STR_LOCAL) set(ARCH_CHECK -1 PARENT_SCOPE) diff --git a/catch/unit/memory/hipArray3DCreate.cc b/catch/unit/memory/hipArray3DCreate.cc index a5d3a4bca..ccd5b9b17 100644 --- a/catch/unit/memory/hipArray3DCreate.cc +++ b/catch/unit/memory/hipArray3DCreate.cc @@ -26,7 +26,7 @@ THE SOFTWARE. namespace { void checkArrayIsExpected(const hipArray_t array, const HIP_ARRAY3D_DESCRIPTOR& expected_desc) { // hipArray3DGetDescriptor doesn't currently exist (EXSWCPHIPT-87) -#if HT_AMD +#if HT_AMD || HT_SPIRV std::ignore = array; std::ignore = expected_desc; #else @@ -58,7 +58,7 @@ TEMPLATE_TEST_CASE("Unit_hipArray3DCreate_happy", "", char, uchar2, uint2, int4, HIP_ARRAY3D_DESCRIPTOR desc{}; desc.Format = vec_info::format; desc.NumChannels = vec_info::size; -#if HT_AMD +#if HT_AMD || HT_SPIRV desc.Flags = 0; #else desc.Flags = GENERATE(0, hipArraySurfaceLoadStore, hipArrayTextureGather); @@ -99,7 +99,7 @@ TEMPLATE_TEST_CASE("Unit_hipArray3DCreate_MaxTexture", "", int, uint4, short, us HIP_ARRAY3D_DESCRIPTOR desc{}; desc.Format = vec_info::format; desc.NumChannels = vec_info::size; -#if HT_AMD +#if HT_AMD || HT_SPIRV desc.Flags = 0; #else desc.Flags = GENERATE(0, hipArraySurfaceLoadStore); @@ -337,7 +337,7 @@ TEMPLATE_TEST_CASE("Unit_hipArray3DCreate_Negative_Non2DTextureGather", "", char float2, float4) { CHECK_IMAGE_SUPPORT -#if HT_AMD +#if HT_AMD || HT_SPIRV HipTest::HIP_SKIP_TEST("Texture Gather arrays not supported using AMD backend"); return; #endif diff --git a/catch/unit/memory/hipArrayCommon.hh b/catch/unit/memory/hipArrayCommon.hh index 4d4170060..e99ba43ae 100644 --- a/catch/unit/memory/hipArrayCommon.hh +++ b/catch/unit/memory/hipArrayCommon.hh @@ -45,7 +45,7 @@ __global__ void readFromTexture(T* output, hipTextureObject_t texObj, size_t wid } else { const float v = y / (float)height; if (textureGather) { - output[y * width + x] = tex2Dgather(texObj, u, v, ChannelToRead); + // output[y * width + x] = tex2Dgather(texObj, u, v, ChannelToRead); } else { output[y * width + x] = tex2D(texObj, u, v); } diff --git a/catch/unit/memory/hipArrayCreate.cc b/catch/unit/memory/hipArrayCreate.cc index 29686b2ab..e4c9abe80 100644 --- a/catch/unit/memory/hipArrayCreate.cc +++ b/catch/unit/memory/hipArrayCreate.cc @@ -105,7 +105,7 @@ TEST_CASE("Unit_hipArrayCreate_MultiThread") { // Tests ///////////////////////////////////////// -#if HT_AMD +#if HT_AMD || HT_SPIRV constexpr auto NORMALIZED_COORDINATES = HIP_TRSF_NORMALIZED_COORDINATES; constexpr auto READ_AS_INTEGER = HIP_TRSF_READ_AS_INTEGER; #else diff --git a/catch/unit/memory/hipGetSymbolSizeAddress.cc b/catch/unit/memory/hipGetSymbolSizeAddress.cc index 5c011c7e8..82c4538ae 100644 --- a/catch/unit/memory/hipGetSymbolSizeAddress.cc +++ b/catch/unit/memory/hipGetSymbolSizeAddress.cc @@ -76,7 +76,7 @@ static void HipGetSymbolSizeAddressTest(const void* symbol) { ArrayFindIfNot(read_buffer.data(), expected_value, read_buffer.size()); } -#if HT_AMD +#if HT_AMD || HT_SPIRV || HT_SPIRV #define SYMBOL(expr) &HIP_SYMBOL(expr) #else #define SYMBOL(expr) HIP_SYMBOL(expr) @@ -96,7 +96,7 @@ TEST_CASE("Unit_hipGetSymbolSizeAddress_Positive_Basic") { TEST_CASE("Unit_hipGetSymbolAddress_Negative_Parameters") { // Causes a segfault in CUDA -#if HT_AMD +#if HT_AMD || HT_SPIRV SECTION("devPtr == nullptr") { HIP_CHECK_ERROR(hipGetSymbolAddress(nullptr, SYMBOL(int_var)), hipErrorInvalidValue); } @@ -110,7 +110,7 @@ TEST_CASE("Unit_hipGetSymbolAddress_Negative_Parameters") { TEST_CASE("Unit_hipGetSymbolSize_Negative_Parameters") { // Causes a segfault in CUDA -#if HT_AMD +#if HT_AMD || HT_SPIRV SECTION("size == nullptr") { HIP_CHECK_ERROR(hipGetSymbolSize(nullptr, SYMBOL(int_var)), hipErrorInvalidValue); } diff --git a/catch/unit/memory/hipMalloc3DArray.cc b/catch/unit/memory/hipMalloc3DArray.cc index a0741a105..79d3153da 100644 --- a/catch/unit/memory/hipMalloc3DArray.cc +++ b/catch/unit/memory/hipMalloc3DArray.cc @@ -98,7 +98,7 @@ namespace { void checkArrayIsExpected(hipArray_t array, const hipChannelFormatDesc& expected_desc, const hipExtent& expected_extent, const unsigned int expected_flags) { // hipArrayGetInfo doesn't currently exist (EXSWCPHIPT-87) -#if HT_AMD +#if HT_AMD || HT_SPIRV std::ignore = array; std::ignore = expected_desc; std::ignore = expected_extent; @@ -130,7 +130,7 @@ TEMPLATE_TEST_CASE("Unit_hipMalloc3DArray_happy", "", char, uchar2, uint2, int4, hipArray_t array; const auto desc = hipCreateChannelDesc(); -#if HT_AMD +#if HT_AMD || HT_SPIRV const unsigned int flags = hipArrayDefault; #else const unsigned int flags = @@ -161,7 +161,7 @@ TEMPLATE_TEST_CASE("Unit_hipMalloc3DArray_MaxTexture", "", int, uint4, short, us hipArray_t array; const hipChannelFormatDesc desc = hipCreateChannelDesc(); -#if HT_AMD +#if HT_AMD || HT_SPIRV const unsigned int flag = hipArrayDefault; #else const unsigned int flag = GENERATE(hipArrayDefault, hipArraySurfaceLoadStore); @@ -224,7 +224,7 @@ TEMPLATE_TEST_CASE("Unit_hipMalloc3DArray_MaxTexture", "", int, uint4, short, us } -#if HT_AMD +#if HT_AMD || HT_SPIRV constexpr std::array validFlags{hipArrayDefault}; #else constexpr std::array validFlags{ @@ -312,7 +312,7 @@ TEST_CASE("Unit_hipMalloc3DArray_Negative_InvalidFlags") { hipArray_t array; hipChannelFormatDesc desc = hipCreateChannelDesc(); -#if HT_AMD +#if HT_AMD || HT_SPIRV const unsigned int flag = 0xDEADBEEF; #else const unsigned int flag = @@ -438,7 +438,7 @@ TEST_CASE("Unit_hipMalloc3DArray_Negative_NumericLimit") { // texture gather arrays are only allowed to be 2D TEMPLATE_TEST_CASE("Unit_hipMalloc3DArray_Negative_Non2DTextureGather", "", char, uchar2, short4, float2, float4) { -#if HT_AMD +#if HT_AMD || HT_SPIRV HipTest::HIP_SKIP_TEST("Texture Gather arrays not supported using AMD backend"); return; #endif diff --git a/catch/unit/memory/hipMallocArray.cc b/catch/unit/memory/hipMallocArray.cc index 33db8b8b1..688bdf8b4 100644 --- a/catch/unit/memory/hipMallocArray.cc +++ b/catch/unit/memory/hipMallocArray.cc @@ -518,7 +518,7 @@ TEST_CASE("Unit_hipMallocArray_Negative_DifferentChannelSizes") { INFO("format: " << channelFormatString(channelFormat) << ", x bits: " << bitsX << ", y bits: " << bitsY << ", z bits: " << bitsZ << ", w bits: " << bitsW); -#if HT_AMD +#if HT_AMD || HT_SPIRV unsigned int flag = hipArrayDefault; HIP_CHECK_ERROR(hipMallocArray(&arrayPtr, &desc, width, height, flag), hipErrorInvalidValue); #else @@ -569,7 +569,7 @@ TEST_CASE("Unit_hipMallocArray_Negative_BadFlags") { hipArray_t arrayPtr; SECTION("Flags that dont work with 1D") { -#if HT_AMD +#if HT_AMD || HT_SPIRV // * cudaArrayLayered 0x01 - 1 // * cudaArrayCubemap 0x04 - 4 unsigned int flag = @@ -600,7 +600,7 @@ TEMPLATE_TEST_CASE("Unit_hipMallocArray_Negative_8bitFloat", "", float, float2, // pointer to the array in device memory hipArray_t arrayPtr; -#if HT_AMD +#if HT_AMD || HT_SPIRV unsigned int flags = hipArrayDefault; HIP_CHECK_ERROR(hipMallocArray(&arrayPtr, &desc, 1024, 1024, flags), hipErrorInvalidValue); #else @@ -627,7 +627,7 @@ TEST_CASE("Unit_hipMallocArray_Negative_BadNumberOfBits") { hipArray_t arrayPtr; INFO("Number of bits: " << badBits); -#if HT_AMD +#if HT_AMD || HT_SPIRV unsigned int flag = hipArrayDefault; HIP_CHECK_ERROR(hipMallocArray(&arrayPtr, &desc, 1024, 1024, flag), hipErrorInvalidValue); #else @@ -656,7 +656,7 @@ TEST_CASE("Unit_hipMallocArray_Negative_3ChannelElement") { // pointer to the array in device memory hipArray_t arrayPtr; -#if HT_AMD +#if HT_AMD || HT_SPIRV unsigned int flag = hipArrayDefault; HIP_CHECK_ERROR(hipMallocArray(&arrayPtr, &desc, 1024, 1024, flag), hipErrorInvalidValue); #else @@ -685,7 +685,7 @@ TEST_CASE("Unit_hipMallocArray_Negative_ChannelAfterZeroChannel") { INFO("x: " << desc.x << ", y: " << desc.y << ", z: " << desc.z << ", w: " << desc.w); hipArray_t arrayPtr; -#if HT_AMD +#if HT_AMD || HT_SPIRV unsigned int flag = hipArrayDefault; HIP_CHECK_ERROR(hipMallocArray(&arrayPtr, &desc, 1024, 1024, flag), hipErrorInvalidValue); #else @@ -710,7 +710,7 @@ TEST_CASE("Unit_hipMallocArray_Negative_InvalidChannelFormat") { CAPTURE(formatKind); -#if HT_AMD +#if HT_AMD || HT_SPIRV unsigned int flag = hipArrayDefault; HIP_CHECK_ERROR(hipMallocArray(&arrayPtr, &desc, 1024, 1024, flag), hipErrorInvalidValue); #else @@ -728,7 +728,7 @@ TEST_CASE("Unit_hipMallocArray_Negative_NumericLimit") { hipChannelFormatDesc desc = hipCreateChannelDesc(); size_t size = std::numeric_limits::max(); -#if HT_AMD +#if HT_AMD || HT_SPIRV unsigned int flag = hipArrayDefault; #else unsigned int flag = GENERATE(hipArrayDefault, hipArraySurfaceLoadStore, hipArrayTextureGather); diff --git a/catch/unit/memory/hipMallocMipmappedArray.cc b/catch/unit/memory/hipMallocMipmappedArray.cc index 43da90a3e..cc2d18eb2 100644 --- a/catch/unit/memory/hipMallocMipmappedArray.cc +++ b/catch/unit/memory/hipMallocMipmappedArray.cc @@ -105,7 +105,7 @@ void checkMipmappedArrayIsExpected(hipArray_t level_array, const hipExtent& expected_extent, const unsigned int expected_flags) { // hipArrayGetInfo doesn't currently exist (EXSWCPHIPT-87) -#if HT_AMD +#if HT_AMD || HT_SPIRV std::ignore = level_array; std::ignore = expected_desc; std::ignore = expected_extent; @@ -134,7 +134,7 @@ void checkMipmappedArrayIsExpected(hipArray_t level_array, TEMPLATE_TEST_CASE("Unit_hipMallocMipmappedArray_happy", "", char, uint2, int4, short4, float) { hipMipmappedArray_t array; const auto desc = hipCreateChannelDesc(); -#if HT_AMD +#if HT_AMD || HT_SPIRV const unsigned int flags = hipArrayDefault; #else const unsigned int flags = @@ -162,7 +162,7 @@ TEMPLATE_TEST_CASE("Unit_hipMallocMipmappedArray_happy", "", char, uint2, int4, } } -#if HT_AMD +#if HT_AMD || HT_SPIRV constexpr std::array validFlags{hipArrayDefault}; #else constexpr std::array validFlags{ @@ -247,7 +247,7 @@ TEST_CASE("Unit_hipMallocMipmappedArray_Negative_InvalidFlags") { hipMipmappedArray_t array; hipChannelFormatDesc desc = hipCreateChannelDesc(); -#if HT_AMD +#if HT_AMD || HT_SPIRV const unsigned int flag = 0xDEADBEEF; #else const unsigned int flag = @@ -368,7 +368,7 @@ TEST_CASE("Unit_hipMallocMipmappedArray_Negative_NumericLimit") { // texture gather arrays are only allowed to be 2D TEMPLATE_TEST_CASE("Unit_hipMallocMipmappedArray_Negative_Non2DTextureGather", "", char, uchar2, float2) { -#if HT_AMD +#if HT_AMD || HT_SPIRV HipTest::HIP_SKIP_TEST("Texture Gather arrays not supported using AMD backend"); return; #endif diff --git a/catch/unit/memory/hipMemcpyAsync.cc b/catch/unit/memory/hipMemcpyAsync.cc index a57449e84..7e73384c3 100644 --- a/catch/unit/memory/hipMemcpyAsync.cc +++ b/catch/unit/memory/hipMemcpyAsync.cc @@ -58,10 +58,10 @@ TEST_CASE("Unit_hipMemcpyAsync_Positive_Synchronization_Behavior") { false); } - SECTION("Device memory to device Memory No CU") { - MemcpyDtoDSyncBehavior(std::bind(hipMemcpyAsync, _1, _2, _3, hipMemcpyDeviceToDeviceNoCU, - nullptr),false); - } + // SECTION("Device memory to device Memory No CU") { + // MemcpyDtoDSyncBehavior(std::bind(hipMemcpyAsync, _1, _2, _3, hipMemcpyDeviceToDeviceNoCU, + // nullptr),false); + // } SECTION("Host memory to host memory") { MemcpyHtoHSyncBehavior(std::bind(hipMemcpyAsync, _1, _2, _3, hipMemcpyHostToHost, nullptr), diff --git a/catch/unit/memory/hipPointerGetAttribute.cc b/catch/unit/memory/hipPointerGetAttribute.cc index 16101a628..e0e1e3d6e 100644 --- a/catch/unit/memory/hipPointerGetAttribute.cc +++ b/catch/unit/memory/hipPointerGetAttribute.cc @@ -76,7 +76,7 @@ TEST_CASE("Unit_hipPointerGetAttribute_MemoryTypes") { REQUIRE(datatype == hipMemoryTypeDevice); } -#if HT_AMD +#if HT_AMD || HT_SPIRV SECTION("Malloc Array Allocation") { hipArray_t B_d; hipChannelFormatDesc desc = hipCreateChannelDesc(); @@ -205,7 +205,7 @@ TEST_CASE("Unit_hipPointerGetAttribute_BufferID") { hipPointerGetAttribute API with HIP_POINTER_ATTRIBUTE_DEVICE_ORDINAL and ensure that it matches with CUDA result */ -#if HT_AMD +#if HT_AMD || HT_SPIRV TEST_CASE("Unit_hipPointerGetAttribute_HostDeviceOrdinal") { size_t Nbytes = 0; Nbytes = N * sizeof(int); @@ -276,7 +276,7 @@ TEST_CASE("Unit_hipPointerGetAttribute_Negative") { == hipErrorInvalidValue); } SECTION("Pass nullptr to device attribute") { -#if HT_AMD +#if HT_AMD || HT_SPIRV REQUIRE(hipPointerGetAttribute(&data, HIP_POINTER_ATTRIBUTE_DEVICE_POINTER, nullptr) == hipErrorInvalidValue); #else @@ -311,7 +311,7 @@ TEST_CASE("Unit_hipPointerGetAttribute_Negative") { REQUIRE(hipPointerGetAttribute(&data, static_cast(-1), reinterpret_cast(A_h)) == hipErrorInvalidValue); } -#if HT_AMD +#if HT_AMD || HT_SPIRV SECTION("Pass HIP_POINTER_ATTRIBUTE_IS_GPU_DIRECT_RDMA_CAPABLE" "not supported by HIP") { REQUIRE(hipPointerGetAttribute(&data, diff --git a/catch/unit/module/CMakeLists.txt b/catch/unit/module/CMakeLists.txt index beb52bb50..3ba4812f4 100644 --- a/catch/unit/module/CMakeLists.txt +++ b/catch/unit/module/CMakeLists.txt @@ -29,8 +29,8 @@ set(TEST_SRC hipModuleLaunchKernel.cc hipModuleGetGlobal.cc hipModuleGetTexRef.cc - hipModuleLaunchCooperativeKernel.cc - hipModuleLaunchCooperativeKernelMultiDevice.cc + # hipModuleLaunchCooperativeKernel.cc + # hipModuleLaunchCooperativeKernelMultiDevice.cc hipFuncGetAttribute.cc ) diff --git a/catch/unit/printf/CMakeLists.txt b/catch/unit/printf/CMakeLists.txt index 992326b35..e0b064fe1 100644 --- a/catch/unit/printf/CMakeLists.txt +++ b/catch/unit/printf/CMakeLists.txt @@ -4,7 +4,7 @@ set(TEST_SRC printfLength.cc printfSpecifiers.cc printfFlagsNonHost.cc - printfSpecifiersNonHost.cc + # printfSpecifiersNonHost.cc # compiler crash in ::lowerTextureFunctions printfHost.cc ) @@ -51,14 +51,15 @@ endif() # Standalone exes add_executable(printfFlags_exe EXCLUDE_FROM_ALL printfFlags_exe.cc) -add_executable(printfLength_exe EXCLUDE_FROM_ALL printfLength_exe.cc) -add_executable(printfSpecifiers_exe EXCLUDE_FROM_ALL printfSpecifiers_exe.cc) +# add_executable(printfLength_exe EXCLUDE_FROM_ALL printfLength_exe.cc) +# add_executable(printfSpecifiers_exe EXCLUDE_FROM_ALL printfSpecifiers_exe.cc) add_executable(printfFlagsNonHost_exe EXCLUDE_FROM_ALL printfFlagsNonHost_exe.cc) -add_executable(printfSpecifiersNonHost_exe EXCLUDE_FROM_ALL printfSpecifiersNonHost_exe.cc) +# add_executable(printfSpecifiersNonHost_exe EXCLUDE_FROM_ALL printfSpecifiersNonHost_exe.cc) +# Unhandled constant expr: ptr addrspace(4) inttoptr (i64 -1144570581550241922 to ptr addrspace(4)) add_dependencies(build_tests printfFlags_exe) -add_dependencies(build_tests printfLength_exe) -add_dependencies(build_tests printfSpecifiers_exe) +# add_dependencies(build_tests printfLength_exe) +# add_dependencies(build_tests printfSpecifiers_exe) add_dependencies(build_tests printfFlagsNonHost_exe) -add_dependencies(build_tests printfSpecifiersNonHost_exe) +# add_dependencies(build_tests printfSpecifiersNonHost_exe) diff --git a/catch/unit/stream/hipStreamACb_MultiThread.cc b/catch/unit/stream/hipStreamACb_MultiThread.cc index 9d1a780de..5ad263c7a 100644 --- a/catch/unit/stream/hipStreamACb_MultiThread.cc +++ b/catch/unit/stream/hipStreamACb_MultiThread.cc @@ -33,7 +33,7 @@ static std::atomic Cb_count{0}, Data_mismatch{0}; static hipStream_t mystream; static float *A1_h, *C1_h; -#if HT_AMD +#if HT_AMD || HT_SPIRV #define HIPRT_CB #endif diff --git a/catch/unit/stream/hipStreamACb_StrmSyncTiming.cc b/catch/unit/stream/hipStreamACb_StrmSyncTiming.cc index a60760496..b13e3c2f4 100644 --- a/catch/unit/stream/hipStreamACb_StrmSyncTiming.cc +++ b/catch/unit/stream/hipStreamACb_StrmSyncTiming.cc @@ -29,7 +29,7 @@ multiple Threads. #include #include -#ifdef __HIP_PLATFORM_AMD__ +#if HT_AMD || HT_SPIRV #define HIPRT_CB #endif diff --git a/catch/unit/stream/hipStreamAddCallback.cc b/catch/unit/stream/hipStreamAddCallback.cc index 567cfa168..9ee42453f 100644 --- a/catch/unit/stream/hipStreamAddCallback.cc +++ b/catch/unit/stream/hipStreamAddCallback.cc @@ -31,7 +31,7 @@ Testcase Scenarios : #define UNUSED(expr) do { (void)(expr); } while (0) -#ifdef __HIP_PLATFORM_AMD__ +#if HT_AMD || HT_SPIRV #define HIPRT_CB #endif diff --git a/catch/unit/streamperthread/CMakeLists.txt b/catch/unit/streamperthread/CMakeLists.txt index e18594271..5caf8577e 100644 --- a/catch/unit/streamperthread/CMakeLists.txt +++ b/catch/unit/streamperthread/CMakeLists.txt @@ -4,8 +4,8 @@ set(TEST_SRC hipStreamPerThread_Event.cc hipStreamPerThread_MultiThread.cc hipStreamPerThread_DeviceReset.cc - hipStreamPerThrdTsts.cc - hipStreamPerThrdCompilerOptn.cc +# hipStreamPerThrdTsts.cc # cooperative_groups; +# hipStreamPerThrdCompilerOptn.cc # cooperative_groups ) if(HIP_PLATFORM MATCHES "amd") diff --git a/catch/unit/texture/CMakeLists.txt b/catch/unit/texture/CMakeLists.txt index 1fb227814..dcd097f20 100644 --- a/catch/unit/texture/CMakeLists.txt +++ b/catch/unit/texture/CMakeLists.txt @@ -29,8 +29,8 @@ set(TEST_SRC hipTextureObj2D.cc hipSimpleTexture3D.cc hipTextureRef2D.cc - hipSimpleTexture1DLayered.cc - hipSimpleTexture2DLayered.cc + # hipSimpleTexture1DLayered.cc + # hipSimpleTexture2DLayered.cc hipBindTex2DPitch.cc hipBindTexRef1DFetch.cc hipTex1DFetchCheckModes.cc diff --git a/catch/unit/warp/CMakeLists.txt b/catch/unit/warp/CMakeLists.txt index 5fded69b5..f77a6dff7 100644 --- a/catch/unit/warp/CMakeLists.txt +++ b/catch/unit/warp/CMakeLists.txt @@ -1,8 +1,8 @@ # Common Tests - Test independent of all platforms set(TEST_SRC - warp_ballot.cc - warp_any.cc - warp_all.cc + # warp_ballot.cc # coop groups + # warp_any.cc + # warp_all.cc ) if(HIP_PLATFORM MATCHES "amd") From 0f06552177f9e7246d0f21a8d04517563ecbd517 Mon Sep 17 00:00:00 2001 From: Paulius Velesko Date: Mon, 11 Mar 2024 04:29:56 +0200 Subject: [PATCH 07/18] RTC Testing Fixes * Improper string substitutions for kernel locations * Missing includes Still fails to compile - seems like this path is not used? --- catch/CMakeLists.txt | 16 +++++++++++++++- catch/include/hip_test_common.hh | 5 ++++- catch/include/hip_test_rtc.hh | 1 + catch/kernels_path.h.in | 6 ++++++ 4 files changed, 26 insertions(+), 2 deletions(-) create mode 100644 catch/kernels_path.h.in diff --git a/catch/CMakeLists.txt b/catch/CMakeLists.txt index bbbd15b87..4cdad0632 100644 --- a/catch/CMakeLists.txt +++ b/catch/CMakeLists.txt @@ -132,7 +132,21 @@ option(RTC_TESTING "Run tests using HIP RTC to compile the kernels" OFF) if (RTC_TESTING) add_definitions(-DRTC_TESTING=ON) endif() -add_definitions(-DKERNELS_PATH="${CMAKE_CURRENT_SOURCE_DIR}/kernels/") + +# The following does not work +# add_definitions(-DKERNELS_PATH="${CMAKE_CURRENT_SOURCE_DIR}/kernels/") +# In file included from /space/pvelesko/chipStar/hip-tests-update/hip-tests/catch/hipTestMain/main.cc:3: +# In file included from /space/pvelesko/chipStar/hip-tests-update/hip-tests/catch/./include/hip_test_common.hh:37: +# /space/pvelesko/chipStar/hip-tests-update/hip-tests/catch/./include/hip_test_rtc.hh:110:36: error: use of undeclared identifier 'tests' +# :1:68: note: expanded from macro 'KERNELS_PATH' +# 1 | #define KERNELS_PATH /space/pvelesko/chipStar/hip-tests-update/hip-tests/catch/kernels/ +configure_file( + "${CMAKE_CURRENT_SOURCE_DIR}/kernels_path.h.in" + "${CMAKE_CURRENT_BINARY_DIR}/kernels_path.h" +) + +# Include the generated header file directory +include_directories("${CMAKE_CURRENT_BINARY_DIR}") set(CATCH_BUILD_DIR catch_tests) execute_process(COMMAND ${CMAKE_COMMAND} -E diff --git a/catch/include/hip_test_common.hh b/catch/include/hip_test_common.hh index 96e7b27e4..834efe8ef 100644 --- a/catch/include/hip_test_common.hh +++ b/catch/include/hip_test_common.hh @@ -34,6 +34,9 @@ THE SOFTWARE. #include #include #include +// Had to add this include to make the code compile +// error: use of undeclared identifier 'launchRTCKernel' +#include "hip_test_rtc.hh" #define HIP_PRINT_STATUS(status) INFO(hipGetErrorName(status) << " at line: " << __LINE__); @@ -352,7 +355,7 @@ void launchKernel(Kernel kernel, Dim numBlocks, Dim numThreads, std::uint32_t me kernel<<>>(std::forward(args)...); #else launchRTCKernel(kernel, numBlocks, numThreads, memPerBlock, stream, - std::forward(packedArgs)...); + std::forward(args)...); #endif HIP_CHECK(hipGetLastError()); } diff --git a/catch/include/hip_test_rtc.hh b/catch/include/hip_test_rtc.hh index 11ef6a165..1ae277841 100644 --- a/catch/include/hip_test_rtc.hh +++ b/catch/include/hip_test_rtc.hh @@ -34,6 +34,7 @@ THE SOFTWARE. #include #include "hip/hip_runtime_api.h" #include "hip_test_context.hh" +#include "kernels_path.h" namespace HipTest { diff --git a/catch/kernels_path.h.in b/catch/kernels_path.h.in new file mode 100644 index 000000000..867ef6eeb --- /dev/null +++ b/catch/kernels_path.h.in @@ -0,0 +1,6 @@ +#ifndef KERNELS_PATH_H +#define KERNELS_PATH_H + +#define KERNELS_PATH "@CMAKE_CURRENT_SOURCE_DIR@/kernels/" + +#endif \ No newline at end of file From 4e952d145c0a43f5a6c433db0cde6649eb73ea8a Mon Sep 17 00:00:00 2001 From: Paulius Velesko Date: Mon, 11 Mar 2024 13:02:38 +0200 Subject: [PATCH 08/18] WIP --- catch/CMakeLists.txt | 23 +++++-------------- .../external/Catch2/cmake/Catch2/Catch.cmake | 16 ++----------- catch/kernels/CMakeLists.txt | 5 ---- .../performance/stream/hipStreamWaitValue.cc | 2 +- catch/unit/CMakeLists.txt | 18 +++++++++------ catch/unit/dynamicLoading/CMakeLists.txt | 2 ++ catch/unit/kernel/CMakeLists.txt | 6 ++++- catch/unit/memory/CMakeLists.txt | 15 ++++++------ catch/unit/memory/hipGetSymbolSizeAddress.cc | 2 +- catch/unit/memory/hipMemcpyAsync.cc | 8 +++---- catch/unit/module/CMakeLists.txt | 6 ++--- catch/unit/streamperthread/CMakeLists.txt | 4 ++-- .../streamperthread/hipStreamPerThrdTsts.cc | 2 +- catch/unit/texture/CMakeLists.txt | 8 +++++-- 14 files changed, 52 insertions(+), 65 deletions(-) diff --git a/catch/CMakeLists.txt b/catch/CMakeLists.txt index 4cdad0632..88c459310 100644 --- a/catch/CMakeLists.txt +++ b/catch/CMakeLists.txt @@ -206,17 +206,7 @@ message(STATUS "CMAKE HIP ARCHITECTURES: ${CMAKE_HIP_ARCHITECTURES}") # preference to pass arch - # OFFLOAD_ARCH_STR # rocm_agent_enumerator -if(HIP_PLATFORM STREQUAL "spirv" AND NOT MASTER_PROJECT) - # for CHIP-SPV, OFFLOAD_ARCH_STR will be passed to CMAKE_CXX_FLAGS which will make - # compile-only flags (such as -x hip) to be passed to linker. This will cause strange errors. - # These could be removed, but then we will need to add manual linking of libCHIP.so so - # might as well just use a different approach for CHIP-SPV path. - # set(OFFLOAD_ARCH_STR ${HIP_OFFLOAD_COMPILE_OPTIONS_BUILD}) - message(WARNING "HIP_PLATFORM is spirv, OFFLOAD_ARCH_STR is set to ${OFFLOAD_ARCH_STR}") -elseif(HIP_PLATFORM STREQUAL "spirv" AND MASTER_PROJECT) - # TODO skip setting this & enforce use of hipcc - message(FATAL_ERROR "building hip-tests with HIP_PLATFORM=spirv is not supported when building as a standalone project") -elseif(NOT DEFINED OFFLOAD_ARCH_STR +if(NOT DEFINED OFFLOAD_ARCH_STR AND EXISTS "${ROCM_PATH}/bin/rocm_agent_enumerator" AND HIP_PLATFORM STREQUAL "amd" AND UNIX) execute_process(COMMAND "${ROCM_PATH}/bin/rocm_agent_enumerator" @@ -321,12 +311,11 @@ add_subdirectory(unit ${CATCH_BUILD_DIR}/unit) add_subdirectory(ABM ${CATCH_BUILD_DIR}/ABM) add_subdirectory(kernels ${CATCH_BUILD_DIR}/kernels) add_subdirectory(hipTestMain ${CATCH_BUILD_DIR}/hipTestMain) -# add_subdirectory(stress ${CATCH_BUILD_DIR}/stress) -# add_subdirectory(TypeQualifiers ${CATCH_BUILD_DIR}/TypeQualifiers) -# add_subdirectory(perftests ${CATCH_BUILD_DIR}/perftests) -# add_subdirectory(multiproc ${CATCH_BUILD_DIR}/multiproc) -# add_subdirectory(performance ${CATCH_BUILD_DIR}/performance) - +add_subdirectory(stress ${CATCH_BUILD_DIR}/stress) +add_subdirectory(TypeQualifiers ${CATCH_BUILD_DIR}/TypeQualifiers) +add_subdirectory(perftests ${CATCH_BUILD_DIR}/perftests) +add_subdirectory(multiproc ${CATCH_BUILD_DIR}/multiproc) +add_subdirectory(performance ${CATCH_BUILD_DIR}/performance) cmake_policy(POP) diff --git a/catch/external/Catch2/cmake/Catch2/Catch.cmake b/catch/external/Catch2/cmake/Catch2/Catch.cmake index e795b5635..28719ece8 100644 --- a/catch/external/Catch2/cmake/Catch2/Catch.cmake +++ b/catch/external/Catch2/cmake/Catch2/Catch.cmake @@ -284,7 +284,7 @@ function(hip_add_exe_to_target_compile_time_detection) add_executable(${_EXE_NAME} EXCLUDE_FROM_ALL ${SRC_NAME} ${COMMON_SHARED_SRC} $) if(HIP_PLATFORM STREQUAL "amd") target_link_libraries(${_EXE_NAME} hiprtc) - else() + elseif(HIP_PLATFORM STREQUAL "nvidia") target_link_libraries(${_EXE_NAME} nvrtc) endif() endif() @@ -358,8 +358,6 @@ function(hip_add_exe_to_target) get_filename_component(_EXE_NAME ${SRC_NAME} NAME_WLE) endif() - - # Create shared lib of all tests if(NOT RTC_TESTING) add_executable(${_EXE_NAME} EXCLUDE_FROM_ALL ${SRC_NAME} ${COMMON_SHARED_SRC} $ $) @@ -370,8 +368,7 @@ function(hip_add_exe_to_target) elseif(HIP_PLATFORM STREQUAL "nvidia") target_link_libraries(${_EXE_NAME} nvrtc) elseif(HIP_PLATFORM STREQUAL "spirv") - message(FATAL_ERROR "RTC path for SPIRV not yet checked") - target_link_libraries(${_EXE_NAME} spirv) + # nothing extra needed for chipStar else() message(FATAL_ERROR "Unsupported HIP_PLATFORM: ${HIP_PLATFORM}") endif() @@ -397,15 +394,6 @@ function(hip_add_exe_to_target) target_link_libraries(${_EXE_NAME} ${_LINKER_LIBS}) endif() - # link against CHIP-SPV - # Required because CHIP-SPV path is not using OFFLOAD_ARCH_STR - # see hip-tests/CMakeLists.txt - if(NOT MASTER_PROJECT AND HIP_PLATFORM STREQUAL "spirv") - # target_compile_options(${_EXE_NAME} PRIVATE -mllvm -amdgpu-early-inline-all=true) - target_compile_options(${_EXE_NAME} PRIVATE ${HIP_OFFLOAD_COMPILE_OPTIONS_BUILD_}) - target_link_libraries(${_EXE_NAME} CHIP) - endif() - # Add dependency on build_tests to build it on this custom target add_dependencies(${_TEST_TARGET_NAME} ${_EXE_NAME}) diff --git a/catch/kernels/CMakeLists.txt b/catch/kernels/CMakeLists.txt index 8b346dcb2..91e1ab69f 100644 --- a/catch/kernels/CMakeLists.txt +++ b/catch/kernels/CMakeLists.txt @@ -5,9 +5,4 @@ if(NOT RTC_TESTING) add_library(KERNELS EXCLUDE_FROM_ALL OBJECT ${TEST_SRC}) target_compile_options(KERNELS PUBLIC -std=c++17) - - # If compiling as part of CHIP-SPV, add the necessary offload flags which would normally be added by hipcc - if(HIP_PLATFORM STREQUAL "spirv" AND NOT MASTER_PROJECT) - target_compile_options(KERNELS PUBLIC ${HIP_OFFLOAD_COMPILE_OPTIONS_BUILD_}) - endif() endif() diff --git a/catch/performance/stream/hipStreamWaitValue.cc b/catch/performance/stream/hipStreamWaitValue.cc index 5d140d01f..8d4aa3f55 100644 --- a/catch/performance/stream/hipStreamWaitValue.cc +++ b/catch/performance/stream/hipStreamWaitValue.cc @@ -28,7 +28,7 @@ THE SOFTWARE. static int IsStreamWaitValueSupported(int device_id) { int wait_value_supported = 0; -#if HT_AMD +#if HT_AMD || HT_SPIRV HIP_CHECK(hipDeviceGetAttribute(&wait_value_supported, hipDeviceAttributeCanUseStreamWaitValue, device_id)); #else diff --git a/catch/unit/CMakeLists.txt b/catch/unit/CMakeLists.txt index 4fa382c63..8b7c523b7 100644 --- a/catch/unit/CMakeLists.txt +++ b/catch/unit/CMakeLists.txt @@ -25,25 +25,22 @@ add_subdirectory(memory) add_subdirectory(stream_ordered) add_subdirectory(stream) add_subdirectory(event) -# add_subdirectory(occupancy) tex2gather add_subdirectory(device) add_subdirectory(printf) -# add_subdirectory(texture) -# add_subdirectory(surface) add_subdirectory(streamperthread) add_subdirectory(kernel) add_subdirectory(multiThread) add_subdirectory(compiler) add_subdirectory(errorHandling) -# add_subdirectory(cooperativeGrps) +add_subdirectory(cooperativeGrps) add_subdirectory(context) add_subdirectory(device_memory) add_subdirectory(warp) -# add_subdirectory(dynamicLoading) # not supported by chipstar? +add_subdirectory(dynamicLoading) add_subdirectory(g++) -# add_subdirectory(module) +add_subdirectory(module) add_subdirectory(channelDescriptor) -# add_subdirectory(executionControl) # cooperative_groups +add_subdirectory(executionControl) add_subdirectory(math) add_subdirectory(vector_types) add_subdirectory(atomics) @@ -54,6 +51,13 @@ add_subdirectory(syncthreads) add_subdirectory(threadfence) add_subdirectory(virtualMemoryManagement) + +if(NOT HIP_PLATFORM STREQUAL "spirv") + add_subdirectory(occupancy) + add_subdirectory(surface) + add_subdirectory(texture) +endif() + if(HIP_PLATFORM STREQUAL "amd") add_subdirectory(callback) #add_subdirectory(clock) diff --git a/catch/unit/dynamicLoading/CMakeLists.txt b/catch/unit/dynamicLoading/CMakeLists.txt index 57a713370..9b0c99828 100644 --- a/catch/unit/dynamicLoading/CMakeLists.txt +++ b/catch/unit/dynamicLoading/CMakeLists.txt @@ -35,6 +35,8 @@ if(HIP_PLATFORM MATCHES "amd") add_custom_target(libLazyLoad.so COMMAND ${CMAKE_CXX_COMPILER} -fPIC -lpthread -shared ${OFFLOAD_ARCH_STR} ${CMAKE_CURRENT_SOURCE_DIR}/liblazyLoad.cc -I${CMAKE_CURRENT_SOURCE_DIR}/../../include -I${CMAKE_CURRENT_SOURCE_DIR}/../../external/Catch2 -L${HIP_PATH}/${CMAKE_INSTALL_LIBDIR} -o libLazyLoad.so) elseif(HIP_PLATFORM MATCHES "nvidia") add_custom_target(libLazyLoad.so COMMAND ${CMAKE_CXX_COMPILER} -Xcompiler -fPIC -lpthread -shared ${CMAKE_CURRENT_SOURCE_DIR}/liblazyLoad.cc -I${CMAKE_CURRENT_SOURCE_DIR}/../../include -I${CMAKE_CURRENT_SOURCE_DIR}/../../external/Catch2 -I${HIP_PATH}/include/ -o libLazyLoad.so) +elseif(HIP_PLATFORM MATCHES "spirv") +add_custom_target(libLazyLoad.so COMMAND ${CMAKE_CXX_COMPILER} -fPIC -lpthread -shared ${CMAKE_CURRENT_SOURCE_DIR}/liblazyLoad.cc -I${CMAKE_CURRENT_SOURCE_DIR}/../../include -I${CMAKE_CURRENT_SOURCE_DIR}/../../external/Catch2 -I${HIP_PATH}/include/ -o libLazyLoad.so) endif() add_custom_target(bit_extract_kernel.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} ${CMAKE_CURRENT_SOURCE_DIR}/bit_extract_kernel.cpp -o ${CMAKE_CURRENT_BINARY_DIR}/../dynamicLoading/bit_extract_kernel.code -I${HIP_PATH}/include/ -I${CMAKE_CURRENT_SOURCE_DIR}/../../include -L${HIP_PATH}/${CMAKE_INSTALL_LIBDIR}) diff --git a/catch/unit/kernel/CMakeLists.txt b/catch/unit/kernel/CMakeLists.txt index ec6f3997a..b0ae04cd7 100644 --- a/catch/unit/kernel/CMakeLists.txt +++ b/catch/unit/kernel/CMakeLists.txt @@ -22,7 +22,6 @@ set(TEST_SRC hipMemFaultStackAllocation.cc hipLaunchBounds.cc - # hipShflTests.cc hipDynamicShared.cc hipDynamicShared2.cc hipEmptyKernel.cc @@ -40,6 +39,11 @@ if(UNIX) hipPrintfKernel.cc) endif() +if(NOT HIP_PLATFORM MATCHES "spirv") + #error: call to '__shfl' is ambiguous + set(TEST_SRC ${TEST_SRC} hipShflTests.cc) +endif() + # only for AMD if(HIP_PLATFORM MATCHES "amd") set(AMD_SRC diff --git a/catch/unit/memory/CMakeLists.txt b/catch/unit/memory/CMakeLists.txt index 28cb545f7..cdb0f8898 100644 --- a/catch/unit/memory/CMakeLists.txt +++ b/catch/unit/memory/CMakeLists.txt @@ -100,8 +100,8 @@ if(HIP_PLATFORM MATCHES "amd") # For windows build error occurs undefined symbol: hipPointerSetAttribute set(TEST_SRC ${TEST_SRC} hipPointerSetAttribute.cc) endif() -else() - # set(TEST_SRC ${TEST_SRC} hipGetSymbolSizeAddress.cc) +elseif(HIP_PLATFORM MATCHES "nvidia") + set(TEST_SRC ${TEST_SRC} hipGetSymbolSizeAddress.cc) endif() hip_add_exe_to_target(NAME MemoryTest1 @@ -136,8 +136,6 @@ set(TEST_SRC hipMemcpyDtoD.cc hipMemcpyDtoDAsync.cc hipHostMalloc.cc - # /space/pvelesko/chipStar/hip-tests-update/hip-tests/catch/./include/hip_test_common.hh:296:38: error: no matching function for call to 'getExpectedArgs' - # 296 | using expectedArgsTuple = decltype(getExpectedArgs(f)); hipMemcpy_old.cc hipMemcpy_derivatives.cc hipMemcpyAsync.cc @@ -166,7 +164,7 @@ set(TEST_SRC hipMemRangeGetAttributes.cc hipFreeAsync.cc hipMallocAsync.cc - # hipStreamAttachMemAsync.cc # getExpectedArgs + hipStreamAttachMemAsync.cc hipMemRangeGetAttributes_old.cc hipMemGetAddressRange.cc hipMallocMipmappedArray.cc @@ -181,8 +179,11 @@ if(HIP_PLATFORM MATCHES "amd") hipArray3DGetDescriptor.cc) endif() -# set(NOT_FOR_MI200_AND_ABOVE_TEST hipMallocArray.cc hipArrayCreate.cc) # tests not for MI200+ -set(NOT_FOR_MI200_AND_ABOVE_TEST ) # tests not for MI200+ +if (NOT HIP_PLATFORM MATCHES "spirv") + # clang crash on HipTextureLowering.cpp: + # Don't know how to lower this texture use case + set(NOT_FOR_MI200_AND_ABOVE_TEST hipMallocArray.cc hipArrayCreate.cc) # tests not for MI200+ +endif() set(MI200_AND_ABOVE_TARGETS gfx90a gfx940 gfx941 gfx942) function(CheckRejectedArchs OFFLOAD_ARCH_STR_LOCAL) set(ARCH_CHECK -1 PARENT_SCOPE) diff --git a/catch/unit/memory/hipGetSymbolSizeAddress.cc b/catch/unit/memory/hipGetSymbolSizeAddress.cc index 82c4538ae..93c3ac1c6 100644 --- a/catch/unit/memory/hipGetSymbolSizeAddress.cc +++ b/catch/unit/memory/hipGetSymbolSizeAddress.cc @@ -76,7 +76,7 @@ static void HipGetSymbolSizeAddressTest(const void* symbol) { ArrayFindIfNot(read_buffer.data(), expected_value, read_buffer.size()); } -#if HT_AMD || HT_SPIRV || HT_SPIRV +#if HT_AMD || HT_SPIRV #define SYMBOL(expr) &HIP_SYMBOL(expr) #else #define SYMBOL(expr) HIP_SYMBOL(expr) diff --git a/catch/unit/memory/hipMemcpyAsync.cc b/catch/unit/memory/hipMemcpyAsync.cc index 7e73384c3..a57449e84 100644 --- a/catch/unit/memory/hipMemcpyAsync.cc +++ b/catch/unit/memory/hipMemcpyAsync.cc @@ -58,10 +58,10 @@ TEST_CASE("Unit_hipMemcpyAsync_Positive_Synchronization_Behavior") { false); } - // SECTION("Device memory to device Memory No CU") { - // MemcpyDtoDSyncBehavior(std::bind(hipMemcpyAsync, _1, _2, _3, hipMemcpyDeviceToDeviceNoCU, - // nullptr),false); - // } + SECTION("Device memory to device Memory No CU") { + MemcpyDtoDSyncBehavior(std::bind(hipMemcpyAsync, _1, _2, _3, hipMemcpyDeviceToDeviceNoCU, + nullptr),false); + } SECTION("Host memory to host memory") { MemcpyHtoHSyncBehavior(std::bind(hipMemcpyAsync, _1, _2, _3, hipMemcpyHostToHost, nullptr), diff --git a/catch/unit/module/CMakeLists.txt b/catch/unit/module/CMakeLists.txt index 3ba4812f4..abcfbc989 100644 --- a/catch/unit/module/CMakeLists.txt +++ b/catch/unit/module/CMakeLists.txt @@ -29,8 +29,8 @@ set(TEST_SRC hipModuleLaunchKernel.cc hipModuleGetGlobal.cc hipModuleGetTexRef.cc - # hipModuleLaunchCooperativeKernel.cc - # hipModuleLaunchCooperativeKernelMultiDevice.cc + hipModuleLaunchCooperativeKernel.cc + hipModuleLaunchCooperativeKernelMultiDevice.cc hipFuncGetAttribute.cc ) @@ -154,7 +154,7 @@ endif() if(HIP_PLATFORM MATCHES "amd") set(RTCLIB "hiprtc") -else() +elseif(HIP_PLATFORM MATCHES "nvidia") set(RTCLIB "nvrtc") endif() hip_add_exe_to_target(NAME ModuleTest diff --git a/catch/unit/streamperthread/CMakeLists.txt b/catch/unit/streamperthread/CMakeLists.txt index 5caf8577e..e18594271 100644 --- a/catch/unit/streamperthread/CMakeLists.txt +++ b/catch/unit/streamperthread/CMakeLists.txt @@ -4,8 +4,8 @@ set(TEST_SRC hipStreamPerThread_Event.cc hipStreamPerThread_MultiThread.cc hipStreamPerThread_DeviceReset.cc -# hipStreamPerThrdTsts.cc # cooperative_groups; -# hipStreamPerThrdCompilerOptn.cc # cooperative_groups + hipStreamPerThrdTsts.cc + hipStreamPerThrdCompilerOptn.cc ) if(HIP_PLATFORM MATCHES "amd") diff --git a/catch/unit/streamperthread/hipStreamPerThrdTsts.cc b/catch/unit/streamperthread/hipStreamPerThrdTsts.cc index c14b38ce1..a2530d833 100644 --- a/catch/unit/streamperthread/hipStreamPerThrdTsts.cc +++ b/catch/unit/streamperthread/hipStreamPerThrdTsts.cc @@ -50,7 +50,7 @@ THE SOFTWARE. #include "hip/hip_cooperative_groups.h" using namespace std::chrono; using namespace cooperative_groups; -#if HT_AMD +#if HT_AMD || HT_SPIRV #define HIPRT_CB #endif diff --git a/catch/unit/texture/CMakeLists.txt b/catch/unit/texture/CMakeLists.txt index dcd097f20..942208f74 100644 --- a/catch/unit/texture/CMakeLists.txt +++ b/catch/unit/texture/CMakeLists.txt @@ -29,8 +29,6 @@ set(TEST_SRC hipTextureObj2D.cc hipSimpleTexture3D.cc hipTextureRef2D.cc - # hipSimpleTexture1DLayered.cc - # hipSimpleTexture2DLayered.cc hipBindTex2DPitch.cc hipBindTexRef1DFetch.cc hipTex1DFetchCheckModes.cc @@ -53,6 +51,12 @@ set(TEST_SRC hipMipmappedArrayGetLevel.cc ) +if(NOT HIP_PLATFORM MATCHES "spirv") + set(TEST_SRC ${TEST_SRC} + hipSimpleTexture1DLayered.cc + hipSimpleTexture2DLayered.cc) +endif() + # tests not for MI200+ set(NOT_FOR_MI200_AND_ABOVE_TEST tex1Dfetch.cc From 012a386f3ee116c3c2fd58f2aa3950faecd513ec Mon Sep 17 00:00:00 2001 From: Paulius Velesko Date: Wed, 13 Mar 2024 04:33:22 +0200 Subject: [PATCH 09/18] hip_texture_helper - define uchar --- catch/include/hip_texture_helper.hh | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/catch/include/hip_texture_helper.hh b/catch/include/hip_texture_helper.hh index 81dd07e09..39a52095c 100644 --- a/catch/include/hip_texture_helper.hh +++ b/catch/include/hip_texture_helper.hh @@ -1,6 +1,10 @@ #pragma once #include +#ifndef uchar +#define uchar unsigned char +#endif + #define HIP_SAMPLING_VERIFY_EPSILON 0.00001 // The internal precision varies by the GPU family and sometimes within the family. // Thus the following threshold is subject to change. From f33b7d173f7af6a645d6a11a350e3850495c7686 Mon Sep 17 00:00:00 2001 From: Paulius Velesko Date: Wed, 13 Mar 2024 04:33:46 +0200 Subject: [PATCH 10/18] exclude cooperative tests which fail to compile --- catch/unit/cooperativeGrps/CMakeLists.txt | 14 ++++++++++---- 1 file changed, 10 insertions(+), 4 deletions(-) diff --git a/catch/unit/cooperativeGrps/CMakeLists.txt b/catch/unit/cooperativeGrps/CMakeLists.txt index 9732e58f2..54d90d583 100644 --- a/catch/unit/cooperativeGrps/CMakeLists.txt +++ b/catch/unit/cooperativeGrps/CMakeLists.txt @@ -2,7 +2,6 @@ set(TEST_SRC thread_block.cc thread_block_tile.cc - coalesced_group_tiled_partition.cc hipCGThreadBlockType_old.cc hipCGMultiGridGroupType_old.cc hipCGGridGroupType_old.cc @@ -12,13 +11,20 @@ set(TEST_SRC hipLaunchCooperativeKernel_old.cc hipLaunchCooperativeKernelMultiDevice_old.cc multi_grid_group.cc - coalesced_groups_shfl_down_old.cc - coalesced_groups_shfl_up_old.cc hipCGCoalescedGroups_old.cc - coalesced_group.cc grid_group.cc coalesced_tiled_groups_metagrp.cc ) + +# Bugs in SPIRV-LLVM-Translator +if(NOT HIP_PLATFORM STREQUAL "spirv") + set(TEST_SRC ${TEST_SRC} + coalesced_group.cc + coalesced_groups_shfl_down_old.cc + coalesced_groups_shfl_up_old.cc + coalesced_group_tiled_partition.cc) +endif() + if(HIP_PLATFORM STREQUAL "nvidia") set_source_files_properties(hipCGMultiGridGroupType_old.cc PROPERTIES COMPILE_FLAGS "-D_CG_ABI_EXPERIMENTAL -rdc=true -gencode arch=compute_60,code=sm_60 -gencode arch=compute_70,code=sm_70 -gencode arch=compute_80,code=sm_80") set_source_files_properties(hipLaunchCooperativeKernelMultiDevice_old.cc PROPERTIES COMPILE_FLAGS "-D_CG_ABI_EXPERIMENTAL -rdc=true -gencode arch=compute_60,code=sm_60 -gencode arch=compute_70,code=sm_70 -gencode arch=compute_80,code=sm_80") From bb7fd4ec09ecc736aa48c9887f77721b051949b5 Mon Sep 17 00:00:00 2001 From: Paulius Velesko Date: Wed, 13 Mar 2024 08:34:18 +0200 Subject: [PATCH 11/18] only try looking for hipconfig if HIP_VERSION undef --- catch/CMakeLists.txt | 8 +++++--- 1 file changed, 5 insertions(+), 3 deletions(-) diff --git a/catch/CMakeLists.txt b/catch/CMakeLists.txt index 88c459310..bf25558f5 100644 --- a/catch/CMakeLists.txt +++ b/catch/CMakeLists.txt @@ -60,9 +60,11 @@ message(STATUS "ROCM_PATH: ${ROCM_PATH}") set(CMAKE_CXX_COMPILER "${HIP_PATH}/bin/hipcc${EXT}") set(CMAKE_C_COMPILER "${HIP_PATH}/bin/hipcc${EXT}") set(HIPCONFIG_EXECUTABLE "${HIP_PATH}/bin/hipconfig${EXT}") -execute_process(COMMAND ${HIPCONFIG_EXECUTABLE} --version - OUTPUT_VARIABLE HIP_VERSION - OUTPUT_STRIP_TRAILING_WHITESPACE) +if (NOT DEFINED HIP_VERSION) + execute_process(COMMAND ${HIPCONFIG_EXECUTABLE} --version + OUTPUT_VARIABLE HIP_VERSION + OUTPUT_STRIP_TRAILING_WHITESPACE) +endif() # enforce c++17 set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} --std=c++17") From e1768d158711fdeb0858e55f60d8ed1f3fcdec9f Mon Sep 17 00:00:00 2001 From: Paulius Velesko Date: Wed, 13 Mar 2024 08:34:35 +0200 Subject: [PATCH 12/18] add return to suppress warning --- catch/include/hip_test_common.hh | 1 + 1 file changed, 1 insertion(+) diff --git a/catch/include/hip_test_common.hh b/catch/include/hip_test_common.hh index 834efe8ef..848d79221 100644 --- a/catch/include/hip_test_common.hh +++ b/catch/include/hip_test_common.hh @@ -192,6 +192,7 @@ static inline bool IsGfx11() { std::cout << "Have to be either Nvidia or AMD platform, asserting" << std::endl; assert(false); #endif + return false; } From 10a8131ae4890bff9cca7afd2137b5974031f394 Mon Sep 17 00:00:00 2001 From: Paulius Velesko Date: Wed, 13 Mar 2024 10:15:25 +0200 Subject: [PATCH 13/18] Update CatchAddTests.cmake to fix discovery bug --- .../Catch2/cmake/Catch2/CatchAddTests.cmake | 192 +++++++++--------- 1 file changed, 93 insertions(+), 99 deletions(-) diff --git a/catch/external/Catch2/cmake/Catch2/CatchAddTests.cmake b/catch/external/Catch2/cmake/Catch2/CatchAddTests.cmake index 29a8e9aa1..2b13187a7 100644 --- a/catch/external/Catch2/cmake/Catch2/CatchAddTests.cmake +++ b/catch/external/Catch2/cmake/Catch2/CatchAddTests.cmake @@ -29,116 +29,110 @@ function(add_command NAME) set(script "${script}${NAME}(${_args})\n" PARENT_SCOPE) endfunction() +get_filename_component(TEST_EXECUTABLE ${TEST_EXECUTABLE} ABSOLUTE) -foreach(TEST_EXECUTABLE ${TEST_EXE_LIST}) - if(WIN32) - set(TEST_EXECUTABLE ${TEST_EXECUTABLE}.exe) - endif() - get_filename_component(TEST_EXECUTABLE ${TEST_EXECUTABLE} ABSOLUTE) - - # Run test executable to get list of available tests - if(NOT EXISTS "${TEST_EXECUTABLE}") - # exe does not exist moving to the next executable - continue() - endif() - execute_process( - COMMAND ${TEST_EXECUTOR} "${TEST_EXECUTABLE}" ${spec} --list-test-names-only - OUTPUT_VARIABLE output - RESULT_VARIABLE result - WORKING_DIRECTORY "${TEST_WORKING_DIR}" +# Run test executable to get list of available tests +if(NOT EXISTS "${TEST_EXECUTABLE}") + message(FATAL_ERROR + "Specified test executable '${TEST_EXECUTABLE}' does not exist" ) - # Catch --list-test-names-only reports the number of tests, so 0 is... surprising - if(${result} EQUAL 0) - message(WARNING - "Test executable '${TEST_EXECUTABLE}' contains no tests!\n" - ) - elseif(${result} LESS 0) - message(FATAL_ERROR - "Error running test executable '${TEST_EXECUTABLE}':\n" - " Result: ${result}\n" - " Output: ${output}\n" - ) - endif() +endif() +execute_process( + COMMAND ${TEST_EXECUTOR} "${TEST_EXECUTABLE}" ${spec} --list-test-names-only + OUTPUT_VARIABLE output + RESULT_VARIABLE result + WORKING_DIRECTORY "${TEST_WORKING_DIR}" +) +# Catch --list-test-names-only reports the number of tests, so 0 is... surprising +if(${result} EQUAL 0) + message(WARNING + "Test executable '${TEST_EXECUTABLE}' contains no tests!\n" + ) +elseif(${result} LESS 0) + message(FATAL_ERROR + "Error running test executable '${TEST_EXECUTABLE}':\n" + " Result: ${result}\n" + " Output: ${output}\n" + ) +endif() - string(REPLACE "\n" ";" output "${output}") +string(REPLACE "\n" ";" output "${output}") - # Run test executable to get list of available reporters - execute_process( - COMMAND ${TEST_EXECUTOR} "${TEST_EXECUTABLE}" ${spec} --list-reporters - OUTPUT_VARIABLE reporters_output - RESULT_VARIABLE reporters_result - WORKING_DIRECTORY "${TEST_WORKING_DIR}" +# Run test executable to get list of available reporters +execute_process( + COMMAND ${TEST_EXECUTOR} "${TEST_EXECUTABLE}" ${spec} --list-reporters + OUTPUT_VARIABLE reporters_output + RESULT_VARIABLE reporters_result + WORKING_DIRECTORY "${TEST_WORKING_DIR}" +) +if(${reporters_result} EQUAL 0) + message(WARNING + "Test executable '${TEST_EXECUTABLE}' contains no reporters!\n" ) - if(${reporters_result} EQUAL 0) - message(WARNING - "Test executable '${TEST_EXECUTABLE}' contains no reporters!\n" - ) - elseif(${reporters_result} LESS 0) - message(FATAL_ERROR - "Error running test executable '${TEST_EXECUTABLE}':\n" - " Result: ${reporters_result}\n" - " Output: ${reporters_output}\n" - ) - endif() - string(FIND "${reporters_output}" "${reporter}" reporter_is_valid) - if(reporter AND ${reporter_is_valid} EQUAL -1) - message(FATAL_ERROR - "\"${reporter}\" is not a valid reporter!\n" - ) - endif() +elseif(${reporters_result} LESS 0) + message(FATAL_ERROR + "Error running test executable '${TEST_EXECUTABLE}':\n" + " Result: ${reporters_result}\n" + " Output: ${reporters_output}\n" + ) +endif() +string(FIND "${reporters_output}" "${reporter}" reporter_is_valid) +if(reporter AND ${reporter_is_valid} EQUAL -1) + message(FATAL_ERROR + "\"${reporter}\" is not a valid reporter!\n" + ) +endif() - # Prepare reporter - if(reporter) - set(reporter_arg "--reporter ${reporter}") - endif() +# Prepare reporter +if(reporter) + set(reporter_arg "--reporter ${reporter}") +endif() - # Prepare output dir - if(output_dir AND NOT IS_ABSOLUTE ${output_dir}) - set(output_dir "${TEST_WORKING_DIR}/${output_dir}") - if(NOT EXISTS ${output_dir}) - file(MAKE_DIRECTORY ${output_dir}) - endif() +# Prepare output dir +if(output_dir AND NOT IS_ABSOLUTE ${output_dir}) + set(output_dir "${TEST_WORKING_DIR}/${output_dir}") + if(NOT EXISTS ${output_dir}) + file(MAKE_DIRECTORY ${output_dir}) endif() +endif() - # Parse output - foreach(line ${output}) - set(test ${line}) - # Escape characters in test case names that would be parsed by Catch2 - set(test_name ${test}) - foreach(char , [ ]) - string(REPLACE ${char} "\\${char}" test_name ${test_name}) - endforeach(char) - # ...add output dir - if(output_dir) - string(REGEX REPLACE "[^A-Za-z0-9_]" "_" test_name_clean ${test_name}) - set(output_dir_arg "--out ${output_dir}/${output_prefix}${test_name_clean}${output_suffix}") - endif() - - file(RELATIVE_PATH exe_path ${CMAKE_CURRENT_BINARY_DIR} ${TEST_EXECUTABLE}) - - # ...and add to script - add_command(add_test - "${prefix}${test}${suffix}" - ${TEST_EXECUTOR} - "${exe_path}" - "${test_name}" - ${extra_args} - "${reporter_arg}" - "${output_dir_arg}" - ) - add_command(set_tests_properties - "${prefix}${test}${suffix}" - PROPERTIES - ${properties} - ) - list(APPEND tests "${prefix}${test}${suffix}") - endforeach() +# Parse output +foreach(line ${output}) + set(test ${line}) + # Escape characters in test case names that would be parsed by Catch2 + set(test_name ${test}) + foreach(char , [ ]) + string(REPLACE ${char} "\\${char}" test_name ${test_name}) + endforeach(char) + # ...add output dir + if(output_dir) + string(REGEX REPLACE "[^A-Za-z0-9_]" "_" test_name_clean ${test_name}) + set(output_dir_arg "--out ${output_dir}/${output_prefix}${test_name_clean}${output_suffix}") + endif() - # Create a list of all discovered tests, which users may use to e.g. set - # properties on the tests - add_command(set ${TEST_LIST} ${tests}) + file(RELATIVE_PATH exe_path ${CMAKE_CURRENT_BINARY_DIR} ${TEST_EXECUTABLE}) + # ...and add to script + add_command(add_test + "${prefix}${test}${suffix}" + ${TEST_EXECUTOR} + "${exe_path}" + "${test_name}" + ${extra_args} + "${reporter_arg}" + "${output_dir_arg}" + ) + add_command(set_tests_properties + "${prefix}${test}${suffix}" + PROPERTIES + ${properties} + ) + list(APPEND tests "${prefix}${test}${suffix}") endforeach() +# Create a list of all discovered tests, which users may use to e.g. set +# properties on the tests +add_command(set ${TEST_LIST} ${tests}) + # Write CTest script -file(WRITE "${CTEST_FILE}" "${script}") +file(APPEND "${CTEST_FILE}" "${script}") From daaa5ff532bd1a4797103a047ddfacbe6c7e5f15 Mon Sep 17 00:00:00 2001 From: Paulius Velesko Date: Wed, 13 Mar 2024 10:16:09 +0200 Subject: [PATCH 14/18] fix test discovery --- .../external/Catch2/cmake/Catch2/Catch.cmake | 92 ++----------------- 1 file changed, 8 insertions(+), 84 deletions(-) diff --git a/catch/external/Catch2/cmake/Catch2/Catch.cmake b/catch/external/Catch2/cmake/Catch2/Catch.cmake index 28719ece8..d2ce4eb01 100644 --- a/catch/external/Catch2/cmake/Catch2/Catch.cmake +++ b/catch/external/Catch2/cmake/Catch2/Catch.cmake @@ -149,6 +149,7 @@ function(catch_discover_tests_compile_time_detection TARGET TEST_SET) add_custom_command( TARGET ${EXE_NAME} POST_BUILD + BYPRODUCTS "${ctest_tests_file}" COMMAND "${CMAKE_COMMAND}" -D "TEST_TARGET=${EXE_NAME}" -D "TEST_EXECUTABLE=$" @@ -252,88 +253,6 @@ set(_CATCH_DISCOVER_TESTS_SCRIPT CACHE INTERNAL "Catch2 full path to CatchAddTests.cmake helper file" ) - -############################################################################### -# function to be called by all tests -function(hip_add_exe_to_target_compile_time_detection) - set(options) - # NAME EventTest, TEST_SRC src, TEST_TARGET_NAME build_tests - set(args NAME TEST_TARGET_NAME PLATFORM COMPILE_OPTIONS) - set(list_args TEST_SRC LINKER_LIBS COMMON_SHARED_SRC PROPERTY) - cmake_parse_arguments( - PARSE_ARGV 0 - "" # variable prefix - "${options}" - "${args}" - "${list_args}" - ) - - foreach(SRC_NAME ${TEST_SRC}) - if(NOT STANDALONE_TESTS EQUAL "1") - set(_EXE_NAME ${_NAME}) - # take the entire source set for building the executable - set(SRC_NAME ${TEST_SRC}) - else() - # strip extension of src and use exe name as src name - get_filename_component(_EXE_NAME ${SRC_NAME} NAME_WLE) - endif() - - if(NOT RTC_TESTING) - add_executable(${_EXE_NAME} EXCLUDE_FROM_ALL ${SRC_NAME} ${COMMON_SHARED_SRC} $ $) - else () - add_executable(${_EXE_NAME} EXCLUDE_FROM_ALL ${SRC_NAME} ${COMMON_SHARED_SRC} $) - if(HIP_PLATFORM STREQUAL "amd") - target_link_libraries(${_EXE_NAME} hiprtc) - elseif(HIP_PLATFORM STREQUAL "nvidia") - target_link_libraries(${_EXE_NAME} nvrtc) - endif() - endif() - - - - if(UNIX) - set(_LINKER_LIBS ${_LINKER_LIBS} stdc++fs) - set(_LINKER_LIBS ${_LINKER_LIBS} -ldl) - else() - # res files are built resource files using rc files. - # use llvm-rc exe to build the res files - # Thes are used to populate the properties of the built executables - if(EXISTS "${PROP_RC}/catchProp.res") - set(_LINKER_LIBS ${_LINKER_LIBS} "${PROP_RC}/catchProp.res") - endif() - #set(_LINKER_LIBS ${_LINKER_LIBS} -noAutoResponse) - endif() - - if(DEFINED _LINKER_LIBS) - target_link_libraries(${_EXE_NAME} ${_LINKER_LIBS}) - endif() - - # Add dependency on build_tests to build it on this custom target - add_dependencies(${_TEST_TARGET_NAME} ${_EXE_NAME}) - # add_dependencies(${_TEST_TARGET_NAME} ${_EXE_NAME}) - - if (DEFINED _PROPERTY) - set_property(TARGET ${_EXE_NAME} PROPERTY ${_PROPERTY}) - endif() - - if (DEFINED _COMPILE_OPTIONS) - target_compile_options(${_EXE_NAME} PUBLIC ${_COMPILE_OPTIONS}) - endif() - foreach(arg IN LISTS _UNPARSED_ARGUMENTS) - message(WARNING "Unparsed arguments: ${arg}") - endforeach() - get_property(crosscompiling_emulator - TARGET ${_EXE_NAME} - PROPERTY CROSSCOMPILING_EMULATOR - ) - set(_EXE_NAME_LIST ${_EXE_NAME_LIST} ${_EXE_NAME}) - if(NOT STANDALONE_TESTS EQUAL "1") - break() - endif() - endforeach() - catch_discover_tests("${_EXE_NAME_LIST}" "${_NAME}" PROPERTIES SKIP_REGULAR_EXPRESSION "HIP_SKIP_THIS_TEST") -endfunction() - ############################################################################### # current staging # function to be called by all tests @@ -415,6 +334,11 @@ function(hip_add_exe_to_target) endforeach() - catch_discover_tests("${_EXE_NAME_LIST}" "${_NAME}" PROPERTIES SKIP_REGULAR_EXPRESSION "HIP_SKIP_THIS_TEST") -endfunction() + + if(DEFINED CATCH2_DISCOVER_TESTS_COMPILE_TIME AND CATCH2_DISCOVER_TESTS_COMPILE_TIME) + catch_discover_tests_compile_time_detection("${_EXE_NAME_LIST}" "${_NAME}" PROPERTIES SKIP_REGULAR_EXPRESSION "HIP_SKIP_THIS_TEST") + else() + catch_discover_tests("${_EXE_NAME_LIST}" "${_NAME}" PROPERTIES SKIP_REGULAR_EXPRESSION "HIP_SKIP_THIS_TEST") + endif() +endfunction() From b101acc9e3fbeb22ff9c0dd4dca0fc51b15d5e27 Mon Sep 17 00:00:00 2001 From: Paulius Velesko Date: Thu, 14 Mar 2024 09:45:20 +0200 Subject: [PATCH 15/18] more HT_SPIRV --- catch/ABM/AddKernels/add.cc | 2 +- catch/multiproc/hipMemCoherencyTstMProc.cc | 26 +++++++++++----------- catch/unit/memory/hipHostMalloc.cc | 6 ++--- catch/unit/memory/hipHostRegister.cc | 8 +++---- 4 files changed, 21 insertions(+), 21 deletions(-) diff --git a/catch/ABM/AddKernels/add.cc b/catch/ABM/AddKernels/add.cc index 1b7c56cdf..186097b99 100644 --- a/catch/ABM/AddKernels/add.cc +++ b/catch/ABM/AddKernels/add.cc @@ -7,7 +7,7 @@ template __global__ void add(T* a, T* b, T* c, size_t size) { } TEMPLATE_TEST_CASE("ABM_AddKernel_MultiTypeMultiSize", "", int, long, float, long long, double) { - auto size = GENERATE(as{}, 100, 500, 1000); + auto size = GENERATE(as{}, 100, 500); TestType *d_a, *d_b, *d_c; auto res = hipMalloc(&d_a, sizeof(TestType) * size); REQUIRE(res == hipSuccess); diff --git a/catch/multiproc/hipMemCoherencyTstMProc.cc b/catch/multiproc/hipMemCoherencyTstMProc.cc index c80068ca3..b93128c36 100644 --- a/catch/multiproc/hipMemCoherencyTstMProc.cc +++ b/catch/multiproc/hipMemCoherencyTstMProc.cc @@ -52,7 +52,7 @@ __global__ void CoherentTst(int *ptr, int PeakClk) { } __global__ void CoherentTst_gfx11(int *ptr, int PeakClk) { -#if HT_AMD +#if HT_AMD || HT_SPIRV // Incrementing the value by 1 int64_t GpuFrq = int64_t(PeakClk) * 1000; int64_t StrtTck = wall_clock64(); @@ -124,7 +124,7 @@ static void TstCoherency(int *Ptr, bool HmmMem) { /* Test case description: The following test validates if fine grain behavior is observed or not with memory allocated using malloc()*/ // The following test is failing on Nvidia platform hence disabled it for now -#if HT_AMD +#if HT_AMD || HT_SPIRV TEST_CASE("Unit_malloc_CoherentTst") { hipDeviceProp_t prop; HIPCHECK(hipGetDeviceProperties(&prop, 0)); @@ -155,7 +155,7 @@ TEST_CASE("Unit_malloc_CoherentTst") { /* Test case description: The following test validates if coarse grain memory behavior is observed or not with memory allocated using malloc()*/ // The following test is failing on Nvidia platform hence disabling it for now -#if HT_AMD +#if HT_AMD || HT_SPIRV TEST_CASE("Unit_malloc_CoherentTstWthAdvise") { hipDeviceProp_t prop; HIPCHECK(hipGetDeviceProperties(&prop, 0)); @@ -187,7 +187,7 @@ TEST_CASE("Unit_malloc_CoherentTstWthAdvise") { /* Test case description: The following test validates if fine memory behavior is observed or not with memory allocated using mmap()*/ // The following test is failing on Nvidia platform hence disabling it for now -#if HT_AMD +#if HT_AMD || HT_SPIRV TEST_CASE("Unit_mmap_CoherentTst") { hipDeviceProp_t prop; HIPCHECK(hipGetDeviceProperties(&prop, 0)); @@ -224,7 +224,7 @@ TEST_CASE("Unit_mmap_CoherentTst") { /* Test case description: The following test validates if coarse grain memory behavior is observed or not with memory allocated using mmap()*/ // The following test is failing on Nvidia platform hence disabling it for now -#if HT_AMD +#if HT_AMD || HT_SPIRV TEST_CASE("Unit_mmap_CoherentTstWthAdvise") { hipDeviceProp_t prop; HIPCHECK(hipGetDeviceProperties(&prop, 0)); @@ -269,7 +269,7 @@ TEST_CASE("Unit_mmap_CoherentTstWthAdvise") { /* Test Case Description: The following test checks if the memory is accessible when HIP_HOST_COHERENT is set to 0*/ // The following test is AMD specific test hence skipping for Nvidia -#if HT_AMD +#if HT_AMD || HT_SPIRV TEST_CASE("Unit_hipHostMalloc_WthEnv0Flg1") { if ((setenv("HIP_HOST_COHERENT", "0", 1)) != 0) { WARN("Unable to turn on HIP_HOST_COHERENT, hence terminating the Test case!"); @@ -310,7 +310,7 @@ TEST_CASE("Unit_hipHostMalloc_WthEnv0Flg1") { /* Test Case Description: The following test checks if the memory is accessible when HIP_HOST_COHERENT is set to 0*/ // The following test is AMD specific test hence skipping for Nvidia -#if HT_AMD +#if HT_AMD || HT_SPIRV TEST_CASE("Unit_hipHostMalloc_WthEnv0Flg2") { if ((setenv("HIP_HOST_COHERENT", "0", 1)) != 0) { WARN("Unable to turn on HIP_HOST_COHERENT, hence terminating the Test case!"); @@ -351,7 +351,7 @@ TEST_CASE("Unit_hipHostMalloc_WthEnv0Flg2") { /* Test Case Description: The following test checks if the memory is accessible when HIP_HOST_COHERENT is set to 0*/ // The following test is AMD specific test hence skipping for Nvidia -#if HT_AMD +#if HT_AMD || HT_SPIRV TEST_CASE("Unit_hipHostMalloc_WthEnv0Flg3") { if ((setenv("HIP_HOST_COHERENT", "0", 1)) != 0) { WARN("Unable to turn on HIP_HOST_COHERENT, hence terminating the Test case!"); @@ -392,7 +392,7 @@ TEST_CASE("Unit_hipHostMalloc_WthEnv0Flg3") { /* Test Case Description: The following test checks if the memory is accessible when HIP_HOST_COHERENT is set to 0*/ // The following test is AMD specific test hence skipping for Nvidia -#if HT_AMD +#if HT_AMD || HT_SPIRV TEST_CASE("Unit_hipHostMalloc_WthEnv0Flg4") { if ((setenv("HIP_HOST_COHERENT", "0", 1)) != 0) { WARN("Unable to turn on HIP_HOST_COHERENT, hence terminating the Test case!"); @@ -434,7 +434,7 @@ TEST_CASE("Unit_hipHostMalloc_WthEnv0Flg4") { /* Test Case Description: The following test checks if the memory exhibits fine grain behavior when HIP_HOST_COHERENT is set to 1*/ // The following test is AMD specific test hence skipping for Nvidia -#if HT_AMD +#if HT_AMD || HT_SPIRV TEST_CASE("Unit_hipHostMalloc_WthEnv1") { if ((setenv("HIP_HOST_COHERENT", "1", 1)) != 0) { WARN("Unable to turn on HIP_HOST_COHERENT, hence terminating the Test case!"); @@ -473,7 +473,7 @@ TEST_CASE("Unit_hipHostMalloc_WthEnv1") { /* Test Case Description: The following test checks if the memory exhibits fine grain behavior when HIP_HOST_COHERENT is set to 1*/ // The following test is AMD specific test hence skipping for Nvidia -#if HT_AMD +#if HT_AMD || HT_SPIRV TEST_CASE("Unit_hipHostMalloc_WthEnv1Flg1") { if ((setenv("HIP_HOST_COHERENT", "1", 1)) != 0) { WARN("Unable to turn on HIP_HOST_COHERENT, hence terminating the Test case!"); @@ -511,7 +511,7 @@ TEST_CASE("Unit_hipHostMalloc_WthEnv1Flg1") { /* Test Case Description: The following test checks if the memory exhibits fine grain behavior when HIP_HOST_COHERENT is set to 1*/ // The following test is AMD specific test hence skipping for Nvidia -#if HT_AMD +#if HT_AMD || HT_SPIRV TEST_CASE("Unit_hipHostMalloc_WthEnv1Flg2") { if ((setenv("HIP_HOST_COHERENT", "1", 1)) != 0) { WARN("Unable to turn on HIP_HOST_COHERENT, hence terminating the Test case!"); @@ -549,7 +549,7 @@ TEST_CASE("Unit_hipHostMalloc_WthEnv1Flg2") { /* Test Case Description: The following test checks if the memory exhibits fine grain behavior when HIP_HOST_COHERENT is set to 1*/ // The following test is AMD specific test hence skipping for Nvidia -#if HT_AMD +#if HT_AMD || HT_SPIRV TEST_CASE("Unit_hipHostMalloc_WthEnv1Flg3") { if ((setenv("HIP_HOST_COHERENT", "1", 1)) != 0) { WARN("Unable to turn on HIP_HOST_COHERENT, hence terminating the Test case!"); diff --git a/catch/unit/memory/hipHostMalloc.cc b/catch/unit/memory/hipHostMalloc.cc index 6e4c06eb2..5b14fb0cf 100644 --- a/catch/unit/memory/hipHostMalloc.cc +++ b/catch/unit/memory/hipHostMalloc.cc @@ -46,7 +46,7 @@ std::vector syncMsg = {"event", "stream", "device"}; static constexpr int numElements{1024 * 16}; static constexpr size_t sizeBytes{numElements * sizeof(int)}; -#if HT_AMD +#if HT_AMD || HT_SPIRV static __global__ void kerTestMemAccess(char *buf) { size_t myId = threadIdx.x + blockDim.x * blockIdx.x; buf[myId] = VALUE; @@ -162,7 +162,7 @@ This testcase verifies the hipHostMalloc API by passing nullptr to the pointer variable */ TEST_CASE("Unit_hipHostMalloc_Negative") { -#if HT_AMD +#if HT_AMD || HT_SPIRV { // Stimulate error condition: int* A = nullptr; @@ -280,7 +280,7 @@ TEST_CASE("Unit_hipHostMalloc_AllocateMoreThanAvailGPUMemory") { } } -#if HT_AMD +#if HT_AMD || HT_SPIRV TEST_CASE("Unit_hipHostMalloc_AllocateUseMoreThanAvailGPUMemory") { char* A = nullptr; size_t maxGpuMem = 0, availableMem = 0; diff --git a/catch/unit/memory/hipHostRegister.cc b/catch/unit/memory/hipHostRegister.cc index 2110ee928..ab5400c38 100644 --- a/catch/unit/memory/hipHostRegister.cc +++ b/catch/unit/memory/hipHostRegister.cc @@ -45,7 +45,7 @@ static constexpr auto LEN{1024 * 1024}; static constexpr auto LARGE_CHUNK_LEN{100 * LEN}; static constexpr auto SMALL_CHUNK_LEN{10 * LEN}; -#if HT_AMD +#if HT_AMD || HT_SPIRV #define TEST_SKIP(arch, msg) \ if (std::string::npos == arch.find("xnack+")) {\ HipTest::HIP_SKIP_TEST(msg);\ @@ -534,7 +534,7 @@ TEST_CASE("Unit_hipHostRegister_AsyncApis") { HIP_CHECK(hipGetDeviceProperties(&prop, 0)); std::string arch = prop.gcnArchName; bool useRegPtrInDev = false; -#if HT_AMD +#if HT_AMD || HT_SPIRV if (std::string::npos == arch.find("xnack+")) { useRegPtrInDev = false; } else { @@ -594,7 +594,7 @@ TEST_CASE("Unit_hipHostRegister_Graphs") { HIP_CHECK(hipGetDeviceProperties(&prop, 0)); std::string arch = prop.gcnArchName; bool useRegPtrInDev = false; -#if HT_AMD +#if HT_AMD || HT_SPIRV if (std::string::npos == arch.find("xnack+")) { useRegPtrInDev = false; } else { @@ -660,7 +660,7 @@ TEST_CASE("Unit_hipHostRegister_Graphs") { free(B); } -#if HT_AMD +#if HT_AMD || HT_SPIRV /** * Test Description * ------------------------ From 6a57a99660620de9270086942638aacc75497bd1 Mon Sep 17 00:00:00 2001 From: Paulius Velesko Date: Fri, 15 Mar 2024 03:27:12 +0200 Subject: [PATCH 16/18] fix one HT_SPIRV --- catch/include/utils.hh | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/catch/include/utils.hh b/catch/include/utils.hh index 457d215bf..16e44ed66 100644 --- a/catch/include/utils.hh +++ b/catch/include/utils.hh @@ -150,10 +150,10 @@ __global__ void Iota(T* const out, size_t pitch, size_t w, size_t h, size_t d) { inline void LaunchDelayKernel(const std::chrono::milliseconds interval, const hipStream_t stream = nullptr) { int ticks_per_ms = 0; - #if HT_AMD || HT_SPIRV + #if HT_AMD HIPCHECK(hipDeviceGetAttribute(&ticks_per_ms, hipDeviceAttributeWallClockRate, 0)); #endif - #if HT_NVIDIA + #if HT_NVIDIA || HT_SPIRV HIPCHECK(hipDeviceGetAttribute(&ticks_per_ms, hipDeviceAttributeClockRate, 0)); #endif Delay<<<1, 1, 0, stream>>>(interval.count(), ticks_per_ms); From 52a8b06b035b95a602f6d18a8ab8f9ce4e92ef4d Mon Sep 17 00:00:00 2001 From: Paulius Velesko Date: Fri, 15 Mar 2024 03:27:33 +0200 Subject: [PATCH 17/18] remove ALL target which required hipcc --- catch/unit/module/CMakeLists.txt | 9 ++++----- 1 file changed, 4 insertions(+), 5 deletions(-) diff --git a/catch/unit/module/CMakeLists.txt b/catch/unit/module/CMakeLists.txt index abcfbc989..61e64609a 100644 --- a/catch/unit/module/CMakeLists.txt +++ b/catch/unit/module/CMakeLists.txt @@ -39,28 +39,28 @@ add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/get_function_module.code -o get_function_module.code -I${HIP_PATH}/include/ --rocm-path=${ROCM_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_target(get_function_module 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/ --rocm-path=${ROCM_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_target(launch_kernel_module 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/ --rocm-path=${ROCM_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_target(get_global_test_module 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/ --rocm-path=${ROCM_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(get_tex_ref_module DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/get_tex_ref_module.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 @@ -101,7 +101,6 @@ add_custom_target(copiousArgKernel.code -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 From bbbfe89edb2386664395677cff46b59abed2904d Mon Sep 17 00:00:00 2001 From: Paulius Velesko Date: Wed, 3 Apr 2024 02:44:09 +0300 Subject: [PATCH 18/18] re-enable some --genco * device-side mallloc and free, wall_clock placeholders were implemented --- catch/multiproc/CMakeLists.txt | 5 +++++ catch/unit/deviceLib/CMakeLists.txt | 4 +--- 2 files changed, 6 insertions(+), 3 deletions(-) diff --git a/catch/multiproc/CMakeLists.txt b/catch/multiproc/CMakeLists.txt index 448c5bf56..9f3f8bcd9 100644 --- a/catch/multiproc/CMakeLists.txt +++ b/catch/multiproc/CMakeLists.txt @@ -32,6 +32,11 @@ hip_add_exe_to_target(NAME MultiProc TEST_SRC ${TEST_SRC} TEST_TARGET_NAME build_tests LINKER_LIBS hiprtc) +elseif(HIP_PLATFORM MATCHES "spirv") +hip_add_exe_to_target(NAME MultiProc + TEST_SRC ${TEST_SRC} + TEST_TARGET_NAME build_tests + ) endif() if(UNIX) diff --git a/catch/unit/deviceLib/CMakeLists.txt b/catch/unit/deviceLib/CMakeLists.txt index 1d7b746f0..7b6a0c3d6 100644 --- a/catch/unit/deviceLib/CMakeLists.txt +++ b/catch/unit/deviceLib/CMakeLists.txt @@ -189,6 +189,4 @@ elseif(HIP_PLATFORM MATCHES "nvidia" OR HIP_PLATFORM MATCHES "spirv") COMPILE_OPTIONS -Wno-deprecated-declarations) #--Wno-deprecated-declarations unrecognized clang++ endif() -if (NOT HIP_PLATFORM MATCHES "spirv") - add_dependencies(build_tests kerDevAllocMultCO.code kerDevWriteMultCO.code kerDevFreeMultCO.code kerDevAllocSingleKer.code) -endif() +add_dependencies(build_tests kerDevAllocMultCO.code kerDevWriteMultCO.code kerDevFreeMultCO.code kerDevAllocSingleKer.code)