[test-suite] Add HIP Tests to External

[ashi1]

Summary

Adding simple HIP Tests, with a simple saxpy.cpp example.
Share the macros using cmake modules with CUDA.

[ronlieb]

Thanks for putting this together.
if i knew anything about buildbots i would approve it.
but best to defer to an expert.

[tra]

I've added @beanz as the reviewer who has more cmake know-how.

I'm glad to see HIP support being added to the testsuite. I don't know what your plans are for bringing up an LLVM build/test bot for HIP. If you want, I could probably run a compile-only HIP bot on one of the VMs that run CUDA bots.

cmake/modules/Variant.cmake

18

This could use further generalization for ".hip" files.

49

Shared files probably should not rely on variables set somewhere else. Offload should probably be a parameter.
Ditto for Variant* variables. They were OK when everything was in the same file, but less so for something intended for wider use.
They should be specified as a parameter instead.

If we have to keep some parameters as external variables, at the very least such dependencies on external variables must be clearly documented.

[ashi1]

In D99997#2672735, @tra wrote:

I've added @beanz as the reviewer who has more cmake know-how.

I'm glad to see HIP support being added to the testsuite. I don't know what your plans are for bringing up an LLVM build/test bot for HIP. If you want, I could probably run a compile-only HIP bot on one of the VMs that run CUDA bots.

We are planning to add a hip buildbot to test both compilation and execution using various ROCm installations. This allows us to further add tests and experiment with buildbots, and may add OpenMP AMDGPU target.

cmake/modules/Variant.cmake

18

Thanks, adding .hip.

49

Thanks, adding these Variant* variables as parameters, and added some comments about the input and outputs of the macro.

[ashi1]

Revised to Artem's comments, and added empty.hip.

[tra]

LGTM overall.

cmake/modules/Variant.cmake

3

Maybe we should rename the file to GPUTestVariant to reflect its purpose better.

I'd also add a brief description here describing that it provides common functions/macros for External/{CUDA,HIP}.

49

Another thing you may consider is whether it would make sense to switch from macros to functions.
I vaguely recall that when I've implemented this originally, cmake had serious trouble dealing with large number of function calls. With thrust tests enabled, there may be thousands of test targets to be created and with functions it took forever. Perhaps cmake can handle it better these days.

Perhaps we could make the macro a private one and instead provide a function using that macro as the public interface.
Once this patch lands, I can check it on thrust tests. If cmake can handle that, then we could fold the macro into the functions. If not, it will be easy enough to switch back to the macro for my own local tests.

[ashi1]

cmake/modules/Variant.cmake

3

Sounds good to me, I will add it.

49

I think it would be better to introduce functions in a separate commit and have this land with macros first.

Since this method works, should we land as is?

[ashi1]

Address Artem's comments.

[tra]

LGTM.

cmake/modules/Variant.cmake

49

I think it would be better to introduce functions in a separate commit and have this land with macros first.

SGTM.

[JonChesterfield]

Tagging @jdoerfert as HIP CI may be of interest. I wonder if an openmp directory added here could be persuaded to run on nvptx or amdgpu hardware, based on what happens to be in the bot.

[tra]

In D99997#2709907, @JonChesterfield wrote:

Tagging @jdoerfert as HIP CI may be of interest. I wonder if an openmp directory added here could be persuaded to run on nvptx or amdgpu hardware, based on what happens to be in the bot.

SGTM.

What external dependencies would be needed for OMP?

[JonChesterfield]

In D99997#2709961, @tra wrote:

In D99997#2709907, @JonChesterfield wrote:

Tagging @jdoerfert as HIP CI may be of interest. I wonder if an openmp directory added here could be persuaded to run on nvptx or amdgpu hardware, based on what happens to be in the bot.

SGTM.

What external dependencies would be needed for OMP?

Great question! Multiple pieces to the answer so here's a table for the current status.
'hsa' here refers to rocr plus roct which then only needs the linux kernel. Both are easily built via cmake.
'cuda' refers to various pieces of nvidia's toolchain, I'm not clear how the pieces are delimited.

	nvptx device	amdgpu device	nvptx host	amdgpu host
build	none	none	cuda (optional)	hsa (required)
execute	cuda	none	cuda	hsa

The host code also uses a few common linux libraries - libelf, libnuma. In the future amdgpu
openmp execution is expected to depend on math functions from rocm-ocml. Doesn't yet.

Build time dependencies can and probably will be reduced to zero required. Runtime can't be
quite so good because there's a kernel level driver in the stack and llvm doesn't bundle libc.

External/CMakeLists.txt

 add_subdirectory(CUDA)
+add_subdirectory(HIP)
 add_subdirectory(HMMER)
 add_subdirectory(Nurbs)
 add_subdirectory(Povray)
 add_subdirectory(SPEC)
 add_subdirectory(skidmarks10)
 add_subdirectory(sollve_vv)

External/CUDA/CMakeLists.txt

 include(External)
+include(Variant)
 llvm_externals_find(TEST_SUITE_CUDA_ROOT "cuda" "CUDA prerequisites")
 
 set(SUPPORTED_GPU_CUDA_7_0
   sm_20 sm_21
   sm_30 sm_32 sm_35 sm_37
   sm_50 sm_52 sm_53)
 set(SUPPORTED_GPU_CUDA_7_5 ${SUPPORTED_GPU_CUDA_7_0})
 

 set(SUPPORTED_GPU_CUDA_10_0 ${SUPPORTED_GPU_CUDA_9_2}
   sm_75)
 set(SUPPORTED_GPU_CUDA_10_1 ${SUPPORTED_GPU_CUDA_10_0})
 set(SUPPORTED_GPU_CUDA_10_2 ${SUPPORTED_GPU_CUDA_10_1})
 
 set(SUPPORTED_GPU_CUDA_11_0 ${SUPPORTED_GPU_CUDA_10_2}
   sm_80)
 
-# Helper macro to extract version number at the end of the string
-# Input: get_version(Var String)
-#    Where String = /some/string/with/version-x.y.z
-# Output:
-#    Sets Var=x.y.z
-macro(get_version Var Path)
-    string(REGEX MATCH "[0-9]+(\\.[0-9]+)*$" ${Var} ${Path})
-endmacro (get_version)
-
-# Helper function to glob CUDA source files and set LANGUAGE property
-# to CXX on each of them. Sets Var in parent scope to the list of
-# files found.
-macro(cuda_glob Var)
-  file(GLOB FileList ${ARGN})
-  foreach(File IN LISTS FileList)
-    if(${File} MATCHES ".*\.cu$")
-      set_source_files_properties(${File} PROPERTIES LANGUAGE CXX)
-    endif()
-  endforeach()
-  set(${Var} ${FileList})
-endmacro(cuda_glob)
-
-macro(create_one_local_test_f Name FileGlob FilterRegex)
-  if (${VariantSuffix} MATCHES ${FilterRegex})
-    cuda_glob(_sources ${FileGlob})
-    set(_executable ${Name}-${VariantSuffix})
-    set(_executable_path ${CMAKE_CURRENT_BINARY_DIR}/${_executable})
-    llvm_test_run()
-    set(REFERENCE_OUTPUT)
-    # Verify reference output if it exists.
-    if(EXISTS ${CMAKE_CURRENT_SOURCE_DIR}/${Name}.reference_output)
-      set(REFERENCE_OUTPUT ${Name}.reference_output)
-      llvm_test_verify(WORKDIR ${CMAKE_CURRENT_BINARY_DIR}
-        ${FPCMP} %o ${REFERENCE_OUTPUT}-${VariantSuffix}
-      )
-      llvm_test_executable(${_executable} ${_sources})
-      llvm_test_data(${_executable}
-                     DEST_SUFFIX "-${VariantSuffix}"
-                     ${REFERENCE_OUTPUT})
-    else()
-      llvm_test_executable(${_executable} ${_sources})
-    endif()
-    target_compile_options(${_executable} PUBLIC ${VariantCPPFLAGS})
-    if(VariantLibs)
-      target_link_libraries(${_executable} ${VariantLibs})
-    endif()
-    add_dependencies(cuda-tests-simple-${VariantSuffix} ${_executable})
-    # Local tests are presumed to be fast.
-    list(APPEND CUDA_SIMPLE_TEST_TARGETS ${_executable}.test)
-  endif()
-endmacro()
-
-macro(create_one_local_test Name FileGlob)
-  create_one_local_test_f(${Name} ${FileGlob} ".*")
-endmacro()
-
 # Create targets for CUDA tests that are part of the test suite.
 macro(create_local_cuda_tests VariantSuffix)
   create_one_local_test(assert assert.cu)
   create_one_local_test(axpy axpy.cu)
   create_one_local_test(algorithm algorithm.cu)
   create_one_local_test(cmath cmath.cu)
   create_one_local_test(complex complex.cu)
   create_one_local_test(math_h math_h.cu)
function(create_cuda_test_variant Std VariantSuffix)
 
   if(DEFINED THRUST_PATH AND (NOT ${Std} IN_LIST "c++14;c++17;c++20"))
     create_thrust_tests(${VariantSuffix})
   endif()
 
   # Target for CUDA tests that take little time to build and run.
   add_custom_target(check-cuda-simple-${VariantSuffix}
     COMMAND ${TEST_SUITE_LIT} ${TEST_SUITE_LIT_FLAGS}
-            ${CUDA_SIMPLE_TEST_TARGETS}
+            ${VARIANT_SIMPLE_TEST_TARGETS}
     WORKING_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}
     DEPENDS cuda-tests-simple-${VariantSuffix}
     USES_TERMINAL)
   add_dependencies(check-cuda-simple check-cuda-simple-${VariantSuffix})
 endfunction(create_cuda_test_variant)
 
 macro(create_cuda_tests)
   message(STATUS "Checking CUDA prerequisites in ${TEST_SUITE_CUDA_ROOT}")

External/HIP/CMakeLists.txt

+include(External)
+include(Variant)
+llvm_externals_find(TEST_SUITE_HIP_ROOT "hip" "HIP prerequisites")
+
+# Create targets for HIP tests that are part of the test suite.
+macro(create_local_hip_tests VariantSuffix)
+  create_one_local_test(saxpy saxpy.cpp)
+endmacro()
+
+function(create_hip_test VariantSuffix)
+  message(STATUS "Creating HIP test variant ${VariantSuffix}")
+  add_custom_target(hip-tests-simple-${VariantSuffix}
+    COMMENT "Build HIP test variant ${VariantSuffix}")
+
+  set(VariantCPPFLAGS ${_HIP_CPPFLAGS})
+  set(VariantLibs ${_HIP_Libs})
+  list(APPEND LDFLAGS ${_HIP_LDFLAGS})
+
+  create_local_hip_tests(${VariantSuffix})
+  add_dependencies(hip-tests-simple hip-tests-simple-${VariantSuffix})
+
+  add_custom_target(check-hip-simple-${VariantSuffix}
+    COMMAND ${TEST_SUITE_LIT} ${TEST_SUITE_LIT_FLAGS}
+            ${VARIANT_SIMPLE_TEST_TARGETS}
+    WORKING_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}
+    DEPENDS hip-tests-simple-${VariantSuffix}
+    USES_TERMINAL)
+  add_dependencies(check-hip-simple check-hip-simple-${VariantSuffix})
+endfunction(create_hip_test)
+
+macro(create_hip_tests)
+  # Find all rocm installations at Externals/hip/ directory.
+  # For ROCm, the path looks like rocm-4.1.0
+  message(STATUS "Checking HIP prerequisites in ${TEST_SUITE_HIP_ROOT}")
+  file(GLOB RocmVersions ${TEST_SUITE_HIP_ROOT}/rocm-*)
+  list(SORT RocmVersions)
+  foreach(RocmDir IN LISTS RocmVersions)
+    get_version(RocmVersion ${RocmDir})
+    message(STATUS "Found ROCm ${RocmVersion}")
+    list(APPEND ROCM_PATHS ${RocmDir})
+    add_library(amdhip64 SHARED IMPORTED)
+    set_property(TARGET amdhip64 PROPERTY IMPORTED_LOCATION
+      ${RocmDir}/lib/libamdhip64.so)
+  endforeach(RocmDir)
+
+  if(NOT ROCM_PATHS)
+    message(SEND_ERROR
+      "There are no ROCm installations in ${TEST_SUITE_HIP_ROOT}")
+    return()
+  endif()
+
+  add_custom_target(hip-tests-simple
+    COMMENT "Build all simple HIP tests")
+  add_custom_target(check-hip-simple
+    COMMENT "Run all simple HIP tests")
+
+  if(NOT AMDGPU_ARCHS)
+    list(APPEND AMDGPU_ARCHS "gfx906;gfx908")
+  endif()
+
+  foreach(_RocmPath ${ROCM_PATHS})
+    get_version(_RocmVersion ${_RocmPath})
+    set(_HIP_Suffix "hip-${_RocmVersion}")
+    set(_HIP_CPPFLAGS -x hip --hip-device-lib-path=${_RocmPath}/amdgcn/bitcode)
+    list(APPEND _HIP_CPPFLAGS -I${_RocmPath}/include)
+    set(_HIP_LDFLAGS -L${_RocmPath}/lib)
+
+    # Unset these for each iteration of rocm path.
+    set(_ArchFlags)
+    set(_ArchList)
+    foreach(_AMDGPUArch IN LISTS AMDGPU_ARCHS)
+      list(APPEND _ArchFlags --offload-arch=${_AMDGPUArch})
+    endforeach()
+    message(STATUS "Building ${_RocmPath} targets for ${AMDGPU_ARCHS}")
+    list(APPEND _HIP_CPPFLAGS ${_ArchFlags})
+
+    set(_HIP_Libs amdhip64)
+    set(Offload "hip")
+    create_hip_test(${_HIP_Suffix})
+  endforeach()
+
+  add_custom_target(hip-tests-all DEPENDS hip-tests-simple
+    COMMENT "Build all HIP tests.")
+
+  file(COPY lit.local.cfg DESTINATION "${CMAKE_CURRENT_BINARY_DIR}")
+endmacro(create_hip_tests)
+
+if(TEST_SUITE_HIP_ROOT)
+  create_hip_tests()
+endif()

External/HIP/README

+HIP Tests
+==========
+
+HIP tests are enabled if cmake is invoked with
+-DTEST_SUITE_EXTERNALS_DIR=<externals path> and specified externals
+directory contains at least one ROCm installation.
+
+Expected externals directory structure:
+Externals/
+  hip/
+    rocm-X.Y.Z/ -- One or more ROCm installation.
+
+export EXTERNAL_DIR=/your/Externals/path
+export AMDGPU_ARCHS=gfx906;gfx908 # List of AMDGPU archs to compile
+export CLANG_DIR=/your/clang/build/dir
+export TEST_SUITE_DIR=/path/to/test-suite-sources
+
+Configure, build and run tests:
+
+```
+$ mkdir test-suite-build-dir
+$ cd test-suite-build-dir
+$ PATH=$CLANG_DIR/bin:$PATH CXX=clang++ CC=clang cmake -G Ninja -DTEST_SUITE_EXTERNALS_DIR=$EXTERNAL_DIR -DAMDGPU_ARCHS=$AMDGPU_ARCHS -DCMAKE_CXX_COMPILER="$CLANG_DIR/bin/clang++" -DCMAKE_C_COMPILER="$CLANG_DIR/bin/clang" $TEST_SUITE_DIR
+$ ninja hip-tests-simple
+$ ninja check-hip-simple
+```
+
+This will build every test for each of the installed ROCm in the
+$EXTERNAL_DIR/hip location, and run them against the expected
+reference_output.

External/HIP/lit.local.cfg

+# -*- python -*-
+
+import os
+
+hip_env_vars = [
+    'HIP_VISIBLE_DEVICES',
+    'LD_LIBRARY_PATH',
+]
+
+for var in hip_env_vars:
+  if var in os.environ:
+    config.environment[var] = os.environ[var]
+
+config.traditional_output = True

External/HIP/saxpy.cpp

+#include <iostream>
+
+#include <hip/hip_runtime.h>
+
+#define N  (1024 * 500)
+
+__global__ void saxpy(float a, float* x, float* y) {
+    size_t tid = blockIdx.x * blockDim.x + threadIdx.x;
+    if (tid < N) y[tid] = a * x[tid] + y[tid];
+}
+
+int main() {
+
+  const float a = 100.0f;
+  float* x = (float*)malloc(N * sizeof(float));
+  float* y = (float*)malloc(N * sizeof(float));
+
+  // Initialize the input data.
+  for (size_t i = 0; i < N; ++i) {
+    x[i] = static_cast<float>(i);
+    y[i] = static_cast<float>(i * 2);
+  }
+
+  // Make a copy for the GPU implementation.
+  float* d_x;
+  float* d_y;
+  hipMalloc((void**)&d_x, N * sizeof(float));
+  hipMalloc((void**)&d_y, N * sizeof(float));
+  hipMemcpy(d_x, x, N * sizeof(float), hipMemcpyHostToDevice);
+  hipMemcpy(d_y, y, N * sizeof(float), hipMemcpyHostToDevice);
+
+  // CPU implementation of saxpy.
+  for (int i = 0; i < N; i++) {
+    y[i] = a * x[i] + y[i];
+  }
+
+  // Launch a GPU kernel to compute the saxpy.
+  saxpy<<<(N+255)/256, 256>>>(a, d_x, d_y);
+
+  // Copy the device results to host.
+  float* h_y = (float*)malloc(N * sizeof(float));
+  hipDeviceSynchronize();
+  hipMemcpy(h_y, d_y, N * sizeof(float), hipMemcpyDeviceToHost);
+
+  // Verify the results match CPU.
+  int errors = 0;
+  for (int i = 0; i < N; i++) {
+    if (fabs(y[i] - h_y[i]) > fabs(y[i] * 0.0001f))
+      errors++;
+  }
+  if (errors != 0)
+    std::cout << errors << " errors" << std::endl;
+  else
+    std::cout << "PASSED!" << std::endl;
+
+  free(h_y);
+  free(x);
+  free(y);
+  hipFree(d_x);
+  hipFree(d_y);
+  return errors;
+}

External/HIP/saxpy.reference_output

+PASSED!
+exit 0

cmake/modules/Variant.cmake

+include(External)
+
+# Helper macro to extract version number at the end of the string
+# Input: get_version(Var String)
+#    Where String = /some/string/with/version-x.y.z
+# Output:
+#    Sets Var=x.y.z
+macro(get_version Var Path)
+    string(REGEX MATCH "[0-9]+(\\.[0-9]+)*$" ${Var} ${Path})
+endmacro (get_version)
+
+# Helper function to glob CUDA source files and set LANGUAGE property
+# to CXX on each of them. Sets Var in parent scope to the list of
+# files found.
+macro(cuda_glob Var)
+  file(GLOB FileList ${ARGN})
+  foreach(File IN LISTS FileList)
+    if(${File} MATCHES ".*\.cu$")
+      set_source_files_properties(${File} PROPERTIES LANGUAGE CXX)
+    endif()
+  endforeach()
+  set(${Var} ${FileList})
+endmacro(cuda_glob)
+
+macro(create_one_local_test_f Name FileGlob FilterRegex)
+  if (${VariantSuffix} MATCHES ${FilterRegex})
+    cuda_glob(_sources ${FileGlob})
+    set(_executable ${Name}-${VariantSuffix})
+    set(_executable_path ${CMAKE_CURRENT_BINARY_DIR}/${_executable})
+    llvm_test_run()
+    set(REFERENCE_OUTPUT)
+    # Verify reference output if it exists.
+    if(EXISTS ${CMAKE_CURRENT_SOURCE_DIR}/${Name}.reference_output)
+      set(REFERENCE_OUTPUT ${Name}.reference_output)
+      llvm_test_verify(WORKDIR ${CMAKE_CURRENT_BINARY_DIR}
+        ${FPCMP} %o ${REFERENCE_OUTPUT}-${VariantSuffix}
+      )
+      llvm_test_executable(${_executable} ${_sources})
+      llvm_test_data(${_executable}
+                     DEST_SUFFIX "-${VariantSuffix}"
+                     ${REFERENCE_OUTPUT})
+    else()
+      llvm_test_executable(${_executable} ${_sources})
+    endif()
+    target_compile_options(${_executable} PUBLIC ${VariantCPPFLAGS})
+    if(VariantLibs)
+      target_link_libraries(${_executable} ${VariantLibs})
+    endif()
+    if (${Offload} MATCHES "hip")
+      add_dependencies(hip-tests-simple-${VariantSuffix} ${_executable})
+    else()
+      add_dependencies(cuda-tests-simple-${VariantSuffix} ${_executable})
+    endif()
+    # Local tests are presumed to be fast.
+    list(APPEND VARIANT_SIMPLE_TEST_TARGETS ${_executable}.test)
+  endif()
+endmacro()
+
+macro(create_one_local_test Name FileGlob)
+  create_one_local_test_f(${Name} ${FileGlob} ".*")
+endmacro()
+