[CUDA][HIP] add test for uses of std::array on device code

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)

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 <stdio.h>
+
+#if __cplusplus >= 201103L
+
+#include <array>
+#include <assert.h>
+
+#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<int, 4>
+          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>>>();
+  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

External/CUDA/array.reference_output

@@ -0,0 +1,2 @@
+Success!
+exit 0

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)