diff --git a/catch/unit/compiler/CMakeLists.txt b/catch/unit/compiler/CMakeLists.txt index c0ad0131d1..1f4626df8c 100644 --- a/catch/unit/compiler/CMakeLists.txt +++ b/catch/unit/compiler/CMakeLists.txt @@ -19,7 +19,7 @@ if(HIP_PLATFORM MATCHES "amd") TEST_TARGET_NAME build_tests) set(OFFLOAD_ARCH_GENERIC_STR "--offload-arch=gfx9-generic --offload-arch=gfx9-4-generic:sramecc+:xnack- --offload-arch=gfx9-4-generic:sramecc-:xnack- --offload-arch=gfx9-4-generic:xnack+ --offload-arch=gfx10-1-generic --offload-arch=gfx10-3-generic --offload-arch=gfx11-generic --offload-arch=gfx12-generic") - + set(DISABLE_GENERIC_TARGET_ONLY) # Build hipSquareGenericTargetOnly to cover generic targets only @@ -40,7 +40,7 @@ if(HIP_PLATFORM MATCHES "amd") set_property(GLOBAL APPEND PROPERTY G_INSTALL_CUSTOM_TARGETS ${CMAKE_CURRENT_BINARY_DIR}/${GENERIC_TARGET_ONLY_EXE}) set_property(GLOBAL APPEND PROPERTY G_INSTALL_CUSTOM_TARGETS ${CMAKE_CURRENT_BINARY_DIR}/${GENERIC_TARGET_ONLY_COMPRESSED_EXE}) else() - set(DISABLE_GENERIC_TARGET_ONLY "-DNO_GENERIC_TARGET_ONLY_TEST") + set(DISABLE_GENERIC_TARGET_ONLY "-DNO_GENERIC_TARGET_ONLY_TEST") endif() # Build hipSquareGenericTarget to cover generic targets and the specific target @@ -72,4 +72,16 @@ if(HIP_PLATFORM MATCHES "amd") add_dependencies(hipSquareGenericTarget hipSquareGenericTargetOnly) add_dependencies(hipSquareGenericTarget hipSquareGenericTargetOnlyCompressed) endif() + + add_custom_target(hipSpirvTest ALL + COMMAND ${CMAKE_CXX_COMPILER} ${CMAKE_CURRENT_SOURCE_DIR}/hipSpirvTest.cc + ${CMAKE_CURRENT_SOURCE_DIR}/../../hipTestMain/hip_test_context.cc + ${CMAKE_CURRENT_SOURCE_DIR}/../../hipTestMain/main.cc + -I${CMAKE_CURRENT_SOURCE_DIR}/../../include + -I${CMAKE_CURRENT_SOURCE_DIR}/../../external/Catch2 + -I${CMAKE_CURRENT_SOURCE_DIR}/../../external/picojson + -I${HIP_PATH}/include/ --rocm-path=${ROCM_PATH} --offload-arch=amdgcnspirv + -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/compiler/hipSpirvTest) + add_dependencies(CompilerTest hipSpirvTest) + endif() diff --git a/catch/unit/compiler/hipSpirvTest.cc b/catch/unit/compiler/hipSpirvTest.cc new file mode 100644 index 0000000000..cd54350fd3 --- /dev/null +++ b/catch/unit/compiler/hipSpirvTest.cc @@ -0,0 +1,28 @@ +/* +Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include + +__global__ void kernel() { asm volatile("v_nop" ::: "memory"); } + +// This test case compiles with --offload-arch=amdgcnspirv to verify SPIRV mode +TEST_CASE("Unit_test_spirv_mode") { kernel<<<1, 32>>>(); } \ No newline at end of file