[test-suite] Added CMake files for external CUDA tests.

[tra]

Summary

Creates test executables and RUN-only test scripts for all combinations of:

• CUDA version: all versions found in externals/cuda
• C++ dialect: C++98,C++11
• C++ library: libc++, libstdc++ (for all gcc versions found in externals/cuda)

Supported tests:

• simple CUDA executables (included)
• thrust tests (enabled if if thrust sources are found in externals/cuda)

[MatzeB]

This should not affect users without users without cuda stuff installed so this LGTM.

I don't have any experience from the cuda side and do not plan to use this. So consider the rest as a suggestion I got while glancing over the code (but feel free to do whatever you want in the cuda parts).

External/CUDA/CMakeLists.txt

83–88

It's unusual to do stuff like setting CPPFLAGS in a function rather than toplevel. At least in the sense that it is inconsistent with the style in the rest of the test-suite.

114

if(LARGE_PROBLEM_SIZE) is probably enough.

132–159

This looks like magic to discover the host system. Maybe this would better fit in some global configuration variables combines with cmake/cache files for typical configurations? Esp. the hardcoding of the libstdc++ paths looks out of place here...

175–189

ditto.

202–222

ditto.

External/CUDA/README

4–6

Note that as a convenience you can even do without this flag: If you checkout your externals in test-suite/test-suite-externals it will get picked up automatically at cmake time.

[tra]

External/CUDA/CMakeLists.txt

83–88

Part of the problem is that cmake wants to have one configuration for *everything* while I want to build multiple variants of the same test without creating separate build directory for each, configure things there, etc. because I'd need *a lot* of them. My current test set involves 12 build variants. Once CUDA-8 is out, it 's likely to become 24.

Also, within each build variant I need to build more than one application, each of them potentially requiring its own set of compiler flags. Hence I create thrust app inside a function which gives each application its own set of flags within build variant.

Granted, I may be forcing CMake to do something it was not intended for. I'm new to CMake, so if you have suggestions how to do it better, I'd appreciate your input.

I guess I could move thrust into its own variant iteration loop and set compiler flags that are identical for all variants outside of the loop. It would not do much to change the fact that I'll have large number of targets that each needs its own set of compiler options and libraries to link with.

114

Fixed.

132–159

Not quite. I need to test with different libc++ variants provided in the externals directory. In order to do that I need to provide explicit path to libstdc++ shared library so executable is linked with appropriate -Wl,rpath and picks correct shared library when it's executed. Otherwise I end up loading whatever libstdc++ is installed on host or failing to run because it's not found.

Again, it's probably something I can move out of per-variant loop, but at the moment I can't think of a way to do just let CMake magically pick it only for specific subset of targets I'm creating.

175–189

Similar story here. Per CUDA version installed I need to link with specific cuda runtime library and provide specific include paths. Can't set it globally for everyone. I can move it out of the per-variant loop and do it once, though.

202–222

I need to find *specific* libc++/libstdc++ library to link with. CMake's standard "find me something" facilities are not flexible enough for what I need here.

External/CUDA/README

4–6

Good to know. I'll add this info.

[tra]

I've moved as much as I could from inner loop to outer ones.
Switched most of helper functions to macros as it makes it easier to update state variables in current scope.
Instead of creating variant build targets in separate function's scope,
pass custom options as parameters to llvm_multisource() (see D20221)

[tra]

Updated cmake scripts to use D20842.

Use llvm_test_executable(), retrieve target name it creates, and customize the way we compile/link it using standard cmake functions.

External/CMakeLists.txt

if(NOT ${PATHVAR} AND
     set(${PATHVAR} "${CMAKE_SOURCE_DIR}/test-suite-externals/${NAME}")
   endif()
   if(${PATHVAR})
     message(STATUS "Found ${DESCRIPTION}: ${${PATHVAR}}")
   endif()
 endmacro()
 
 llvm_add_subdirectories(
+  CUDA
   HMMER
   Nurbs
   Povray
   SPEC
   skidmarks10
 )

External/CUDA/CMakeLists.txt

+llvm_externals_find(TEST_SUITE_CUDA_ROOT "cuda" "CUDA prerequisites")
+
+# Helper function 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 in parent scope
+function(get_version Var Path)
+    string(REGEX REPLACE ".*-" "" Version ${Path})
+    set(${Var} ${Version} PARENT_SCOPE)
+endfunction(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.
+function(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} PARENT_SCOPE)
+endfunction(cuda_glob)
+
+function(create_one_local_test Name FileGlob)
+  set(PROG ${Name}-${VariantSuffix})
+  cuda_glob(Source ${FileGlob})
+  llvm_test_run()
+  llvm_multisource()
+  target_link_libraries(${PROG} ${VariantLibs})
+  list(APPEND VARIANT_CUDA_TESTS ${PROG})
+endfunction()
+
+# Create targets for CUDA tests that are part of the test suite.
+function(create_local_cuda_tests VariantSuffix)
+  create_one_local_test(axpy axpy.cu)
+  create_one_local_test(empty empty.cu)
+  set(VARIANT_CUDA_TESTS ${VARIANT_CUDA_TESTS} PARENT_SCOPE)
+endfunction()
+
+macro(thrust_make_test_name TestName TestSourcePath)
+  string(REPLACE "${THRUST_PATH}/testing/" "" _tmp ${TestSourcePath})
+  string(REPLACE "/" "-" _tmp ${_tmp})
+  string(REPLACE "<" "_" _tmp ${_tmp})
+  string(REPLACE ">" "_" _tmp ${_tmp})
+  string(REGEX REPLACE "\.(cpp|cu)$" "" _tmp ${_tmp})
+  set(${TestName} ${_tmp})
+endmacro()
+
+function(create_one_thrust_test_executable TestSourcePath)
+  thrust_make_test_name(TestName ${TestSourcePath})
+  set(PROG thrust-${TestName}-${VariantSuffix})
+  set(Source ${TestSourcePath})
+  llvm_multisource()
+  target_link_libraries(${PROG} ${VariantLibs})
+  set(THRUST_TEST_EXECUTABLE ${PROG} PARENT_SCOPE)
+endfunction()
+
+function(create_one_thrust_test TestSourcePath)
+  thrust_make_test_name(TestName ${TestSourcePath})
+  set(PROG thrust-${TestName}-${VariantSuffix})
+  set(Source ${TestSourcePath})
+  llvm_multisource()
+  target_link_libraries(${PROG} ${VariantLibs})
+  list(APPEND VARIANT_CUDA_TESTS ${PROG})
+  set(VARIANT_CUDA_TESTS ${VARIANT_CUDA_TESTS} PARENT_SCOPE)
+endfunction()
+
+function(create_thrust_test_file Executable TestName)
+  thrust_make_test_name(TestName ${TestName})
+  message(STATUS "Created test ${Executable}-${TestName}")
+  set(executable_path ${CMAKE_CURRENT_BINARY_DIR}/${Executable})
+  if(DEFINED LARGE_PROBLEM_SIZE)
+    llvm_test_run(--sizes=large ${TestName})
+  else()
+    llvm_test_run(${TestName})
+  endif()
+  llvm_add_test("${executable_path}-${TestName}.test" "${executable_path}")
+endfunction()
+
+function(create_thrust_tests VariantSuffix)
+  list(APPEND CPPFLAGS
+    -DTHRUST_HOST_SYSTEM=THRUST_HOST_SYSTEM_CPP
+    -DTHRUST_DEVICE_SYSTEM=THRUST_DEVICE_SYSTEM_CUDA)
+  list(APPEND CPPFLAGS
+    -I${CudaPath}/include
+    -I${THRUST_PATH} -I${THRUST_PATH}/testing)
+
+  # test framework is common for all tests, so we build it once as a
+  # library. Test suite's cmake files don't handle libraries well, so
+  # we need to set compile flags manually.
+  add_library(ThrustTestFrameworkLib-${VariantSuffix}
+    STATIC ${ThrustTestFramework})
+  append_compile_flags(ThrustTestFrameworkLib-${VariantSuffix} ${CPPFLAGS})
+  add_dependencies(ThrustTestFrameworkLib-${VariantSuffix}
+    timeit-host fpcmp-host)
+  list(APPEND VariantLibs ThrustTestFrameworkLib-${VariantSuffix})
+
+  if(DEFINED THRUST_SPLIT_TESTS AND "${THRUST_SPLIT_TESTS}")
+    # Create individual test executable per test source file. This
+    # stresses cmake -- it consumes tons of memory and takes forever
+    # to finish.
+    foreach(TestSourcePath IN LISTS ThrustAllTestSources)
+      create_one_thrust_test_executable(${TestSourcePath})
+      list(APPEND VARIANT_CUDA_TESTS ${THRUST_TEST_EXECUTABLE})
+    endforeach()
+  else()
+    # Single monolitic test executable. Alas this stresses linker
+    # during debug build. Final executable may end up being too large
+    # to link.
+    set(PROG thrust-${VariantSuffix})
+    set(Source ${ThrustTestFramework} ${ThrustAllTestSources})
+    if(DEFINED LARGE_PROBLEM_SIZE)
+      llvm_test_run(--verbose --sizes=large)
+    else()
+      llvm_test_run(--verbose)
+    endif()
+    llvm_multisource()
+    target_link_libraries(${PROG} ${VariantLibs})
+    list(APPEND VARIANT_CUDA_TESTS ${PROG})
+    # We can create one test script per thrust test, but running
+    # thrust tests in parallel is bottlenecked by GPU and startup
+    # overhead, so it's actually way slower than running all tests
+    # sequentially.
+  endif()
+  set(VARIANT_CUDA_TESTS ${VARIANT_CUDA_TESTS} PARENT_SCOPE)
+endfunction()
+
+# Create set of tests for a given {CUDA,C++ standard,C++ library} tuple.
+# Sets VARIANT_CUDA_TESTS targets in parent's scope.
+function(create_cuda_test_variant CudaPath Std GccPath)
+  get_version(CudaVersion ${CudaPath})
+  set(VariantSuffix "cuda-${CudaVersion}")
+  list(APPEND CPPFLAGS --cuda-path=${CudaPath})
+  list(APPEND VariantLibs cudart-${CudaVersion})
+
+  if(Std)
+    set(VariantSuffix "${VariantSuffix}-${Std}")
+    list(APPEND CPPFLAGS -std=${Std})
+    list(APPEND LDFLAGS -std=${Std})
+  endif()
+
+  if(GccPath)
+    get_version(GccVersion ${GccPath})
+    set(VariantSuffix "${VariantSuffix}-libstdc++-${GccVersion}")
+
+    # Tell clang to use libstdc++ and where to find it.
+    list(APPEND CPPFLAGS -gcc-toolchain ${GccPath}/usr/local -stdlib=libstdc++)
+    # Add libstdc++ as link dependency.
+    list(APPEND VariantLibs libstdcxx-${GccVersion})
+  else()
+    # Tell clang to use libc++
+    # We also need to add compiler's include path for cxxabi.h
+    get_filename_component(_compiler_path ${CMAKE_CXX_COMPILER} DIRECTORY)
+    list(APPEND CPPFLAGS -stdlib=libc++ -I${_compiler_path}/../include)
+    list(APPEND LDFLAGS -stdlib=libc++)
+    list(APPEND VariantLibs libcxx)
+  endif()
+
+  message(STATUS "Creating CUDA test variant ${VariantSuffix}")
+  create_local_cuda_tests(${VariantSuffix})
+  if(EXISTS ${THRUST_PATH})
+    create_thrust_tests(${VariantSuffix})
+  endif()
+
+  # Create target to build all tests for this variant
+  add_custom_target(cuda-tests-${VariantSuffix} DEPENDS ${VARIANT_CUDA_TESTS}
+    COMMENT "Build CUDA test variant ${VariantSuffix}")
+  # And pass it up to the caller
+  set(VARIANT_CUDA_TESTS cuda-tests-${VariantSuffix} PARENT_SCOPE)
+endfunction(create_cuda_test_variant)
+
+macro(create_cuda_tests)
+  message(STATUS "Checking CUDA prerequisites in ${TEST_SUITE_CUDA_ROOT}")
+  file(GLOB CudaVersions ${TEST_SUITE_CUDA_ROOT}/cuda-*)
+  foreach(CudaDir IN LISTS CudaVersions)
+    get_version(CudaVersion ${CudaDir})
+    message(STATUS "Found CUDA ${CudaVersion} ${CudaDir}")
+    list(APPEND CUDA_PATHS ${CudaDir})
+    add_library(cudart-${CudaVersion} SHARED IMPORTED)
+    set_property(TARGET cudart-${CudaVersion} PROPERTY IMPORTED_LOCATION
+      ${CudaDir}/lib64/libcudart.so)
+  endforeach(CudaDir)
+
+  if(NOT CUDA_PATHS)
+    message(SEND_ERROR
+      "There are no CUDA installations in ${TEST_SUITE_CUDA_ROOT}")
+    return()
+  endif()
+
+  # set default GPU arch
+  if(NOT CUDA_GPU_ARCH)
+    list(APPEND CUDA_GPU_ARCH sm_35)
+  endif()
+
+  foreach(GpuArch IN LISTS CUDA_GPU_ARCH)
+    list(APPEND CPPFLAGS --cuda-gpu-arch=${GpuArch})
+  endforeach()
+
+  file(GLOB GccVersions ${TEST_SUITE_CUDA_ROOT}/gcc-*)
+  foreach(GccDir IN LISTS GccVersions)
+    get_version(GccVersion ${GccDir})
+    message(STATUS "Found GCC ${GccVersion}")
+    list(APPEND GCC_PATHS ${GccDir})
+    add_library(libstdcxx-${GccVersion} SHARED IMPORTED)
+    set_property(TARGET libstdcxx-${GccVersion} PROPERTY IMPORTED_LOCATION
+      ${GccDir}/usr/local/lib64/libstdc++.so)
+  endforeach(GccDir)
+
+  # Find location of libc++
+  execute_process(
+    COMMAND ${CMAKE_CXX_COMPILER} -print-file-name=libc++.so
+    OUTPUT_VARIABLE _path_to_libcxx
+    OUTPUT_STRIP_TRAILING_WHITESPACE)
+
+  if (EXISTS ${_path_to_libcxx})
+    add_library(libcxx SHARED IMPORTED)
+    set_property(TARGET libcxx PROPERTY IMPORTED_LOCATION ${_path_to_libcxx})
+  else()
+    message(ERROR "Can't find libcxx location.")
+    return()
+  endif()
+
+  if(EXISTS "${TEST_SUITE_CUDA_ROOT}/thrust")
+    set(THRUST_PATH "${TEST_SUITE_CUDA_ROOT}/thrust" CACHE
+      PATH "Thrust library path")
+    message(STATUS "Found Thrust ${THRUST_PATH}")
+    cuda_glob(ThrustTestCppSources ${THRUST_PATH}/testing/*.cpp)
+    cuda_glob(ThrustTestCudaSources ${THRUST_PATH}/testing/*.cu)
+    cuda_glob(ThrustTestCudaBackendSources
+      ${THRUST_PATH}/testing/backend/decompose.cu
+      ${THRUST_PATH}/testing/backend/cuda/*.cu)
+
+    list(APPEND ThrustAllTestSources ${ThrustTestCppSources}
+      ${ThrustTestCudaSources} ${ThrustTestCudaBackendSources})
+    list(APPEND ThrustTestFramework
+      ${THRUST_PATH}/testing/testframework.cpp
+      ${THRUST_PATH}/testing/reduce_large.cu
+      ${THRUST_PATH}/testing/unittest_tester.cu
+      ${THRUST_PATH}/testing/backend/cuda/testframework.cu)
+
+    # Remove test framework files from the list of test files.
+    foreach(File IN LISTS ThrustTestFramework)
+      list(REMOVE_ITEM ThrustAllTestSources ${File})
+    endforeach()
+  endif()
+
+  foreach(CudaPath ${CUDA_PATHS})
+    foreach(Std IN ITEMS "" "c++11")
+      foreach(GccPath IN LISTS GCC_PATHS ITEMS "")
+  set(VARIANT_CUDA_TESTS)
+  create_cuda_test_variant("${CudaPath}" "${Std}" "${GccPath}")
+  list(APPEND ALL_CUDA_TESTS ${VARIANT_CUDA_TESTS})
+      endforeach()
+    endforeach()
+  endforeach()
+
+  # convenience target to build all CUDA tests.
+  add_custom_target(cuda-tests-all DEPENDS ${ALL_CUDA_TESTS}
+    COMMENT "Build all CUDA tests.")
+endmacro(create_cuda_tests)
+
+if(TEST_SUITE_CUDA_ROOT)
+  create_cuda_tests()
+endif()
+

External/CUDA/README

+CUDA Tests
+==========
+
+Cuda tests are enabled if cmake is invoked with
+-DTEST_SUITE_EXTERNALS_DIR=<externals path> and specified externals
+directory contains at least one CUDA installation.
+
+Expected externals directory structure:
+Externals/
+  cuda/
+    cuda-X.Y/  -- One or more CUDA installation.
+    gcc-X.Y.Z/ -- One or more GCC installation for libstdc++.
+
+export EXTERNALS=/your/externals/path
+export CUDA_EXTERNALS=$EXTERNALS/cuda
+export CLANG_DIR=/your/clang/build/dir
+export TEST_SUITE_DIR=/path/to/test-suite-sources
+
+* Externals installation
+  For each cuda version:
+    # bash cuda-linux64-rel-XXXXXX.run -prefix=$CUDA_EXTERNALS/cuda-7.5 -noprompt -nosymlink
+
+  For each GCC version:
+    * extract GCC sources and cd into gcc-X.Y.Z
+    # ../configure --enable-languages=c,c++ --disable-libsanitizer
+    # make bootstrap
+    # make DESTDIR=$CUDA_EXTERNALS/gcc-X.Y.Z install
+
+  Fetch thrust:
+    # cd $CUDA_EXTERNALS
+    # git clone https://github.com/thrust/thrust.git
+
+* 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=$EXTERNALS $TEST_SUITE_DIR 
+  # ninja cuda-all
+  # lit -j1 Externals/CUDA

External/CUDA/axpy.cu

+#include <iostream>
+
+__global__ void axpy(float a, float* x, float* y) {
+  y[threadIdx.x] = a * x[threadIdx.x];
+}
+
+int main(int argc, char* argv[]) {
+  const int kDataLen = 4;
+
+  float a = 2.0f;
+  float host_x[kDataLen] = {1.0f, 2.0f, 3.0f, 4.0f};
+  float host_y[kDataLen];
+
+  // Copy input data to device.
+  float* device_x;
+  float* device_y;
+  cudaMalloc(&device_x, kDataLen * sizeof(float));
+  cudaMalloc(&device_y, kDataLen * sizeof(float));
+  cudaMemcpy(device_x, host_x, kDataLen * sizeof(float),
+             cudaMemcpyHostToDevice);
+
+  // Launch the kernel.
+  axpy<<<1, kDataLen>>>(a, device_x, device_y);
+
+  // Copy output data to host.
+  cudaDeviceSynchronize();
+  cudaMemcpy(host_y, device_y, kDataLen * sizeof(float),
+             cudaMemcpyDeviceToHost);
+
+  // Print the results.
+  for (int i = 0; i < kDataLen; ++i) {
+    std::cout << "y[" << i << "] = " << host_y[i] << "\n";
+  }
+
+  cudaDeviceReset();
+  return 0;
+}

External/CUDA/empty.cu

+int main(int argc, char **argv) { return 0; }