diff --git a/External/CUDA/CMakeLists.txt b/External/CUDA/CMakeLists.txt index 3bc74ecb91..91b8ccc826 100644 --- a/External/CUDA/CMakeLists.txt +++ b/External/CUDA/CMakeLists.txt @@ -63,6 +63,7 @@ macro(create_local_cuda_tests VariantSuffix) list(APPEND CUDA_LOCAL_TESTS assert) list(APPEND CUDA_LOCAL_TESTS axpy) list(APPEND CUDA_LOCAL_TESTS algorithm) + list(APPEND CUDA_LOCAL_TESTS array) list(APPEND CUDA_LOCAL_TESTS cmath) list(APPEND CUDA_LOCAL_TESTS complex) list(APPEND CUDA_LOCAL_TESTS math_h) diff --git a/External/CUDA/array.cu b/External/CUDA/array.cu new file mode 100644 index 0000000000..80196b728b --- /dev/null +++ b/External/CUDA/array.cu @@ -0,0 +1,85 @@ +// Check that we can use std::array on device code +// +// After libstdc++ 15, some internal asserts rely on function that are neither +// constexpr nor device. This can trigger errors when using std::array members +// on device code. +// +// This workaround is implemented in bits/c++config.h + +#include + +#if __cplusplus >= 201103L + +#include +#include + +#if __cplusplus >= 201402L && STDLIB_VERSION >= 2014 +// call the function in a constexpr and a non-constexpr context +#define TEST(expr) \ + do { \ + size_t M = expr; \ + (void)(M); \ + constexpr size_t N = expr; \ + (void)(N); \ + } while (0) +#define MAYBE_CONSTEXPR constexpr +#else +#define TEST(expr) \ + do { \ + size_t M = expr; \ + (void)(M); \ + } while (0) +#define MAYBE_CONSTEXPR +#endif + +MAYBE_CONSTEXPR __host__ __device__ size_t test_array() { + std::array A = {0, 1, 2, 3}; + std::array CtorCopy(A); + std::array CtorMove(std::move(CtorCopy)); + std::array CopyCopy = A; + std::array CopyMove = std::move(CopyCopy); + (void)CtorMove; + (void)CopyMove; + + size_t N = A.size(); + assert(N == 4); + +#if __cplusplus >= 201402L && STDLIB_VERSION >= 2014 + const int fst = A[0]; + assert(fst == 0); +#endif + +#if __cplusplus >= 201703L && STDLIB_VERSION >= 2017 + A[0] = 4; + int snd = A[0]; + assert(snd == 4); +#endif + return N; +} + +__host__ __device__ void do_all_tests() { TEST(test_array()); } + +__global__ void kernel() { do_all_tests(); } + +int main() { + kernel<<<32, 32>>>(); + cudaError_t err = cudaDeviceSynchronize(); + if (err != cudaSuccess) { + printf("CUDA error %d\n", (int)err); + return 1; + } + + do_all_tests(); + + printf("Success!\n"); + return 0; +} + +#else + +int main() { + printf("Success!\n"); + return 0; +} + +#endif diff --git a/External/CUDA/array.reference_output b/External/CUDA/array.reference_output new file mode 100644 index 0000000000..5de1437b90 --- /dev/null +++ b/External/CUDA/array.reference_output @@ -0,0 +1,2 @@ +Success! +exit 0 diff --git a/External/HIP/CMakeLists.txt b/External/HIP/CMakeLists.txt index 11d92fd127..0f74ec883d 100644 --- a/External/HIP/CMakeLists.txt +++ b/External/HIP/CMakeLists.txt @@ -7,6 +7,54 @@ message(STATUS "TEST_SUITE_HIP_ROOT: ${TEST_SUITE_HIP_ROOT}") get_filename_component(HIP_CLANG_PATH ${CMAKE_CXX_COMPILER} DIRECTORY) message(STATUS "HIP_CLANG_PATH: ${HIP_CLANG_PATH}") +# Inspired from create_one_local_test. Runs hipify on the TestSource and then compiles it. +# Search for the reference files next to TestSource. +macro(create_one_hipify_cuda_test TestName TestSource VairantOffload VariantSuffix VariantCPPFlags VariantLibs) + set(_cuda_src "${TestSource}") + set(_hip_src "${TestName}.hip") + set(_hipify_target "${TestName}-hipify") + + set_source_files_properties(${_hip_src} PROPERTIES LANGUAGE CXX) + add_custom_command(OUTPUT ${_hip_src} + COMMAND ${HIPIFY_PERL_EXE} "${_cuda_src}" -o "${_hip_src}" + DEPENDS "${_cuda_src}") + add_custom_target(${_hipify_target} DEPENDS ${_hip_src}) + + set(_executable ${TestName}-${VariantSuffix}) + set(_executable_path ${CMAKE_CURRENT_BINARY_DIR}/${_executable}) + llvm_test_run() + + get_filename_component(_test_source_dir "${TestSource}" DIRECTORY) + get_filename_component(_test_source_name "${TestSource}" NAME_WE) + set(REFERENCE_OUTPUT "${_test_source_dir}/${test_source_name}.reference_output") + if(EXISTS "${REFERENCE_OUTPUT}") + llvm_test_verify(WORKDIR %S + %b/${FPCMP} %o ${REFERENCE_OUTPUT}-${VariantSuffix} + ) + llvm_test_executable(${_executable} ${_hip_src}) + llvm_test_data(${_executable} + DEST_SUFFIX "-${VariantSuffix}" + ${REFERENCE_OUTPUT}) + else() + llvm_test_executable(${_executable} ${_hip_src}) + endif() + + target_compile_options(${_executable} PUBLIC ${VariantCPPFLAGS}) + + # In External/CUDA, tests define a STDLIB_VERSION that matches the C++ + # standard supported by the standard library. + # For the HIP case, we set a huge number and assume that the latest C++ + # standard version is supported by the library. + target_compile_definitions(${_executable} PRIVATE STDLIB_VERSION=9999) + add_dependencies(${_executable} ${_hipify_target}) + if(VariantLibs) + target_link_libraries(${_executable} ${VariantLibs}) + endif() + + add_dependencies(hip-tests-simple-${VariantSuffix} ${_executable}) + list(APPEND VARIANT_SIMPLE_TEST_TARGETS ${_executable}.test) +endmacro() + # Create targets for HIP tests that are part of the test suite. macro(create_local_hip_tests VariantSuffix) set(VariantOffload "hip") @@ -48,6 +96,28 @@ macro(create_local_hip_tests VariantSuffix) "${VariantCPPFLAGS}" "${VariantLibs}") endforeach() + list(APPEND CUDA_LOCAL_TESTS algorithm) + list(APPEND CUDA_LOCAL_TESTS array) + list(APPEND CUDA_LOCAL_TESTS cmath) + list(APPEND CUDA_LOCAL_TESTS complex) + list(APPEND CUDA_LOCAL_TESTS math_h) + list(APPEND CUDA_LOCAL_TESTS new) + + find_program(HIPIFY_PERL_EXE + NAME hipify-perl + PATHS ${_RocmPath}/bin) + + if(HIPIFY_PERL_EXE) + foreach(_cuda_test IN LISTS CUDA_LOCAL_TESTS) + set(_cuda_src "${CMAKE_CURRENT_SOURCE_DIR}/../CUDA/${_cuda_test}.cu") + create_one_hipify_cuda_test(${_cuda_test} ${_cuda_src} + ${VariantOffload} ${VariantSuffix} + "${VariantCPPFLAGS}" "${VariantLibs}") + endforeach() + else() + message(WARNING "hipify-perl not found for ROCm installation in ${_RocmPath}.") + endif() + # Add test for Blender. configure_file(workload/blender/test_blender.sh.in ${CMAKE_CURRENT_BINARY_DIR}/test_blender.sh @ONLY) configure_file(workload/blender/verify_blender.sh.in ${CMAKE_CURRENT_BINARY_DIR}/verify_blender.sh @ONLY)