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/HIP/CMakeLists.txt b/External/HIP/CMakeLists.txt index 3674436513..0420417616 100644 --- a/External/HIP/CMakeLists.txt +++ b/External/HIP/CMakeLists.txt @@ -65,6 +65,7 @@ macro(create_local_hip_tests VariantSuffix) #set_source_files_properties(split-kernel-args.hip PROPERTIES # COMPILE_FLAGS "-mllvm -amdgpu-enable-split-kernel-args") # Add HIP tests to be added to hip-tests-simple + list(APPEND HIP_LOCAL_TESTS array) list(APPEND HIP_LOCAL_TESTS empty) list(APPEND HIP_LOCAL_TESTS with-fopenmp) list(APPEND HIP_LOCAL_TESTS saxpy) diff --git a/External/HIP/array.hip b/External/HIP/array.hip new file mode 100644 index 0000000000..eacb938752 --- /dev/null +++ b/External/HIP/array.hip @@ -0,0 +1,86 @@ +#include "hip/hip_runtime.h" +// 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() { + // Before C++17 only "operator[] const" is constexpr (thus available on + // device). +#if __cplusplus < 201703L && STDLIB_VERSION < 2017 + const +#endif + std::array + A = {0, 1, 2, 3}; + + size_t N = A.size(); + assert(N == 4); + +#if __cplusplus >= 201402L && STDLIB_VERSION >= 2014 + 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>>>(); + hipError_t err = hipDeviceSynchronize(); + if (err != hipSuccess) { + 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