[libomptarget-nvptx] Add testing infrastructure

[Hahnfeld]

Summary

This patch also introduces testing for libomptarget-nvptx
which has been missing until now. I propose to add tests for
all bugs that are fixed in the future.
The target check-libomptarget-nvptx is not run by default because

• we can't determine if there is a GPU plugged into the system.
• it will require the latest Clang compiler. Keeping compatibility with older releases would prevent testing newer code generation developed in trunk.

[Hahnfeld]

libomptarget/deviceRTLs/nvptx/src/parallel.cu

214–223 ↗

(On Diff #164048)

I'm not really sure why this is needed. It was added in https://github.com/clang-ykt/openmp/commit/588e8cbc751299c9f40067dc4d36812e8e49d2bd#diff-6fa63679541dd3fa12162559c1aeac42R279 and I can't find a technical explanation. Because I can't test on Volta (only Pascal) I've copied the code for now, but it would be great to know the background and add an explanation.

[Hahnfeld]

Split testing infrastructure from bugfix.

[grokos]

Great job, adding testing infrastructure has been on our list for a long time!

I don't have any comments, just waiting for the --libomptarget-nvptx-path flag to be accepted into the clang driver. Then we can have a look at the rest of this stack.

[gtbercea]

Considering your comment in the description about requiring latest Clang perhaps you should revisit this patch: D46842

[protze.joachim]

In D51687#1230375, @gtbercea wrote:

Considering your comment in the description about requiring latest Clang perhaps you should revisit this patch: D46842

D46842 only helps if the OpenMP runtime is not built stand-alone.
For the OpenMP runtime tests, we take care that the tests also work with other compatible compilers, and also with older versions.

We might configure different behavior for stand-alone and in-llvm-tree builts.
Considering the necessity of a GPU supporting the features, it makes sense to have a separate target for these tests and executing them manually.

[Hahnfeld]

In D51687#1230375, @gtbercea wrote:

Considering your comment in the description about requiring latest Clang perhaps you should revisit this patch: D46842

In my opinion the two matters are different with regard to an important detail:

1. The linked change in D46842 is about building libomptarget-nvptx with in-tree Clang. As a result the build system would recompile the library whenever there is a change in the Clang sources. For me this drawback outweights the potential gain; in particular you can build the bclib with Clang 7.0 and it will still be inlined when building an application with a later version of Clang (e.g. current trunk), so I don't see the need to use in-tree Clang.
2. This patch adds testing infrastructure which will kick in after libomptarget-nvptx has already been built. At the moment the tests will only work with the latest Clang because it relies on --libomptarget-nvptx-path and code generation that is not yet included in a released version of LLVM. Incidentally I'm listing this as one of the reasons why check-libomptarget-nvptx is not included in neither check-all nor check-openmp. (However I think "we can't determine if there is a GPU plugged into the system" is the one that really matters and we'll probably never be able to solve this, so I don't see enabling these tests by default anyway.)

I hope this clarifies my thoughts that I've brought up in some review threads. From my point of view there is no contradiction to what I propose in this patch and for me the choice of not enabling the tests by default is the logical consequence of my previous statements.

[Hahnfeld]

In D51687#1229696, @grokos wrote:

I don't have any comments, just waiting for the --libomptarget-nvptx-path flag to be accepted into the clang driver. Then we can have a look at the rest of this stack.

Ping, I landed the Clang patch yesterday.

[grokos]

Then this revision is good to go!

openmp/trunk/cmake/OpenMPTesting.cmake

 function(add_openmp_testsuite target comment)
   if (NOT ENABLE_CHECK_TARGETS)
     add_custom_target(${target}
       COMMAND ${CMAKE_COMMAND} -E echo "${target} does nothing, dependencies not found.")
     message(STATUS "${target} does nothing.")
     return()
   endif()
 
-  cmake_parse_arguments(ARG "" "" "DEPENDS" ${ARGN})
+  cmake_parse_arguments(ARG "" "" "DEPENDS;ARGS" ${ARGN})
   # EXCLUDE_FROM_ALL excludes the test ${target} out of check-openmp.
   if (NOT EXCLUDE_FROM_ALL)
     # Register the testsuites and depends for the check-openmp rule.
     set_property(GLOBAL APPEND PROPERTY OPENMP_LIT_TESTSUITES ${ARG_UNPARSED_ARGUMENTS})
     set_property(GLOBAL APPEND PROPERTY OPENMP_LIT_DEPENDS ${ARG_DEPENDS})
   endif()
 
   if (${OPENMP_STANDALONE_BUILD})
+    set(LIT_ARGS ${OPENMP_LIT_ARGS} ${ARG_ARGS})
     add_custom_target(${target}
-      COMMAND ${PYTHON_EXECUTABLE} ${OPENMP_LLVM_LIT_EXECUTABLE} ${OPENMP_LIT_ARGS} ${ARG_UNPARSED_ARGUMENTS}
+      COMMAND ${PYTHON_EXECUTABLE} ${OPENMP_LLVM_LIT_EXECUTABLE} ${LIT_ARGS} ${ARG_UNPARSED_ARGUMENTS}
       COMMENT ${comment}
       DEPENDS ${ARG_DEPENDS}
       ${cmake_3_2_USES_TERMINAL}
     )
   else()
     add_lit_testsuite(${target}
       ${comment}
       ${ARG_UNPARSED_ARGUMENTS}
       DEPENDS clang clang-headers FileCheck ${ARG_DEPENDS}
+      ARGS ${ARG_ARGS}
     )
   endif()
 endfunction()
 
 function(construct_check_openmp_target)
   get_property(OPENMP_LIT_TESTSUITES GLOBAL PROPERTY OPENMP_LIT_TESTSUITES)
   get_property(OPENMP_LIT_DEPENDS GLOBAL PROPERTY OPENMP_LIT_DEPENDS)
 
   # We already added the testsuites themselves, no need to do that again.
   set(EXCLUDE_FROM_ALL True)
   add_openmp_testsuite(check-openmp "Running OpenMP tests" ${OPENMP_LIT_TESTSUITES} DEPENDS ${OPENMP_LIT_DEPENDS})
 endfunction()

openmp/trunk/libomptarget/CMakeLists.txt

if("${CMAKE_SOURCE_DIR}" STREQUAL "${CMAKE_CURRENT_SOURCE_DIR}")
   message(FATAL_ERROR "Direct configuration not supported, please use parent directory!")
 endif()
 
 # Add cmake directory to search for custom cmake functions.
 set(CMAKE_MODULE_PATH ${CMAKE_CURRENT_SOURCE_DIR}/cmake/Modules ${CMAKE_MODULE_PATH})
 
 if(OPENMP_STANDALONE_BUILD)
   # Build all libraries into a common place so that tests can find them.
+  set(CMAKE_ARCHIVE_OUTPUT_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR})
   set(CMAKE_LIBRARY_OUTPUT_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR})
 endif()
 
 # Message utilities.
 include(LibomptargetUtils)
 
 # Get dependencies for the different components of the project.
 include(LibomptargetGetDependencies)
 
 # This is a list of all the targets that are supported/tested right now.
 set (LIBOMPTARGET_ALL_TARGETS "${LIBOMPTARGET_ALL_TARGETS} aarch64-unknown-linux-gnu")
 set (LIBOMPTARGET_ALL_TARGETS "${LIBOMPTARGET_ALL_TARGETS} powerpc64le-ibm-linux-gnu")
 set (LIBOMPTARGET_ALL_TARGETS "${LIBOMPTARGET_ALL_TARGETS} powerpc64-ibm-linux-gnu")
 set (LIBOMPTARGET_ALL_TARGETS "${LIBOMPTARGET_ALL_TARGETS} x86_64-pc-linux-gnu")
 set (LIBOMPTARGET_ALL_TARGETS "${LIBOMPTARGET_ALL_TARGETS} nvptx64-nvidia-cuda")
 
 # Once the plugins for the different targets are validated, they will be added to
 # the list of supported targets in the current system.
 set (LIBOMPTARGET_SYSTEM_TARGETS "")
 
-# Set base directories - required for lit to locate the tests.
-set(LIBOMPTARGET_BASE_DIR ${CMAKE_CURRENT_SOURCE_DIR})
-set(LIBOMPTARGET_BINARY_DIR ${CMAKE_CURRENT_BINARY_DIR})
-
 # If building this library in debug mode, we define a macro to enable
 # dumping progress messages at runtime.
 string( TOLOWER "${CMAKE_BUILD_TYPE}" LIBOMPTARGET_CMAKE_BUILD_TYPE)
 if(LIBOMPTARGET_CMAKE_BUILD_TYPE MATCHES debug)
   add_definitions(-DOMPTARGET_DEBUG)
   add_definitions(-g)
   add_definitions(-O0)
 endif()
 
 include_directories(include)
 
 # Build target agnostic offloading library.
 add_subdirectory(src)
 
 # Retrieve the path to the resulting library so that it can be used for 
 # testing.
 get_target_property(LIBOMPTARGET_LIBRARY_DIR omptarget LIBRARY_OUTPUT_DIRECTORY)
 if(NOT LIBOMPTARGET_LIBRARY_DIR)
   set(LIBOMPTARGET_LIBRARY_DIR ${CMAKE_CURRENT_BINARY_DIR})
 endif()
 
+# Definitions for testing, for reuse when testing libomptarget-nvptx.
+if(OPENMP_STANDALONE_BUILD)
+  set(LIBOMPTARGET_OPENMP_HEADER_FOLDER "${CMAKE_CURRENT_BINARY_DIR}/../runtime/src" CACHE STRING
+    "Path to folder containing omp.h")
+  set(LIBOMPTARGET_OPENMP_HOST_RTL_FOLDER "${CMAKE_CURRENT_BINARY_DIR}/../runtime/src" CACHE STRING
+    "Path to folder containing libomp.so")
+else()
+  set(LIBOMPTARGET_OPENMP_HEADER_FOLDER "${CMAKE_CURRENT_BINARY_DIR}/../runtime/src")
+endif()
+
+
 # Build offloading plugins and device RTLs if they are available.
 add_subdirectory(plugins)
 add_subdirectory(deviceRTLs)
 
 # Add tests.
 add_subdirectory(test)

openmp/trunk/libomptarget/deviceRTLs/nvptx/CMakeLists.txt

if (${LIBOMPTARGET_NVPTX_ENABLE_BCLIB})
 
     # CUDA 9 header files use the nv_weak attribute which clang is not yet prepared
     # to handle. Therefore, we use 'weak' instead. We are compiling only for the
     # device, so it should be equivalent.
     if(CUDA_VERSION_MAJOR GREATER 8)
       set(bc_flags ${bc_flags} -Dnv_weak=weak)
     endif()
 
+    # Create target to build all Bitcode libraries.
+    add_custom_target(omptarget-nvptx-bc)
+
     # Generate a Bitcode library for all the compute capabilities the user requested.
     foreach(sm ${nvptx_sm_list})
       set(cuda_arch --cuda-gpu-arch=sm_${sm})
 
       # Compile CUDA files to bitcode.
       set(bc_files "")
       foreach(src ${cuda_src_files})
         get_filename_component(infile ${src} ABSOLUTE)
foreach(sm ${nvptx_sm_list})
           COMMAND ${LIBOMPTARGET_NVPTX_SELECTED_BC_LINKER}
             -o ${CMAKE_CURRENT_BINARY_DIR}/libomptarget-nvptx-sm_${sm}.bc ${bc_files}
           DEPENDS ${bc_files}
           COMMENT "Linking LLVM bitcode libomptarget-nvptx-sm_${sm}.bc"
       )
       set_property(DIRECTORY APPEND PROPERTY ADDITIONAL_MAKE_CLEAN_FILES libomptarget-nvptx-sm_${sm}.bc)
 
       add_custom_target(omptarget-nvptx-${sm}-bc ALL DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/libomptarget-nvptx-sm_${sm}.bc)
+      add_dependencies(omptarget-nvptx-bc omptarget-nvptx-${sm}-bc)
 
       # Copy library to destination.
       add_custom_command(TARGET omptarget-nvptx-${sm}-bc POST_BUILD
                          COMMAND ${CMAKE_COMMAND} -E copy ${CMAKE_CURRENT_BINARY_DIR}/libomptarget-nvptx-sm_${sm}.bc
                          $<TARGET_FILE_DIR:omptarget-nvptx>)
 
       # Install bitcode library under the lib destination folder.
       install(FILES ${CMAKE_CURRENT_BINARY_DIR}/libomptarget-nvptx-sm_${sm}.bc DESTINATION "${OPENMP_INSTALL_LIBDIR}")
     endforeach()
   endif()
 
+  add_subdirectory(test)
 else()
   libomptarget_say("Not building CUDA offloading device RTL: CUDA tools not found in the system.")
 endif()

openmp/trunk/libomptarget/deviceRTLs/nvptx/test/CMakeLists.txt

+if(NOT OPENMP_TEST_COMPILER_ID STREQUAL "Clang")
+  # Silently return, no need to annoy the user.
+  return()
+endif()
+
+set(deps omptarget-nvptx omptarget omp)
+if(LIBOMPTARGET_NVPTX_ENABLE_BCLIB)
+  set(deps ${deps} omptarget-nvptx-bc)
+endif()
+
+# Don't run by default.
+set(EXCLUDE_FROM_ALL True)
+# Run with only one thread to only launch one application to the GPU at a time.
+add_openmp_testsuite(check-libomptarget-nvptx
+    "Running libomptarget-nvptx tests" ${CMAKE_CURRENT_BINARY_DIR}
+    DEPENDS ${deps} ARGS -j1)
+
+set(LIBOMPTARGET_NVPTX_TEST_FLAGS "" CACHE STRING
+    "Extra compiler flags to send to the test compiler.")
+set(LIBOMPTARGET_NVPTX_TEST_OPENMP_FLAGS
+    "-fopenmp -fopenmp-targets=nvptx64-nvidia-cuda" CACHE STRING
+    "OpenMP compiler flags to use for testing libomptarget-nvptx.")
+
+# Configure the lit.site.cfg.in file
+set(AUTO_GEN_COMMENT "## Autogenerated by libomptarget-nvptx configuration.\n# Do not edit!")
+configure_file(lit.site.cfg.in lit.site.cfg @ONLY)

openmp/trunk/libomptarget/deviceRTLs/nvptx/test/lit.cfg

+# -*- Python -*- vim: set ft=python ts=4 sw=4 expandtab tw=79:
+# Configuration file for the 'lit' test runner.
+
+import os
+import lit.formats
+
+# Tell pylint that we know config and lit_config exist somewhere.
+if 'PYLINT_IMPORT' in os.environ:
+    config = object()
+    lit_config = object()
+
+def prepend_library_path(name, value, sep):
+    if name in config.environment:
+        config.environment[name] = value + sep + config.environment[name]
+    else:
+        config.environment[name] = value
+
+# name: The name of this test suite.
+config.name = 'libomptarget-nvptx'
+
+# suffixes: A list of file extensions to treat as test files.
+config.suffixes = ['.c', '.cpp', '.cc']
+
+# test_source_root: The root path where tests are located.
+config.test_source_root = os.path.dirname(__file__)
+
+# test_exec_root: The root object directory where output is placed
+config.test_exec_root = config.binary_dir
+
+# test format
+config.test_format = lit.formats.ShTest()
+
+# compiler flags
+config.test_flags = " -I " + config.omp_header_directory + \
+    " -L " + config.library_dir + \
+    " --libomptarget-nvptx-path=" + config.library_dir;
+
+if config.omp_host_rtl_directory:
+    config.test_flags = config.test_flags + \
+        " -L " + config.omp_host_rtl_directory
+
+config.test_flags = config.test_flags + " " + config.test_extra_flags
+
+# Setup environment to find dynamic library at runtime.
+prepend_library_path('LD_LIBRARY_PATH', config.library_dir, ":")
+prepend_library_path('LD_LIBRARY_PATH', config.omp_host_rtl_directory, ":")
+
+# Forbid fallback to host.
+config.environment["OMP_TARGET_OFFLOAD"] = "MANDATORY"
+
+# substitutions
+config.substitutions.append(("%compilexx-run-and-check",
+    "%compilexx-and-run | " + config.libomptarget_filecheck + " %s"))
+config.substitutions.append(("%compile-run-and-check",
+    "%compile-and-run | " + config.libomptarget_filecheck + " %s"))
+config.substitutions.append(("%compilexx-and-run", "%compilexx && %run"))
+config.substitutions.append(("%compile-and-run", "%compile && %run"))
+
+config.substitutions.append(("%compilexx",
+    "%clangxx %openmp_flags %flags %s -o %t"))
+config.substitutions.append(("%compile",
+    "%clang %openmp_flags %flags %s -o %t"))
+
+config.substitutions.append(("%clangxx", config.test_cxx_compiler))
+config.substitutions.append(("%clang", config.test_c_compiler))
+config.substitutions.append(("%openmp_flags", config.test_openmp_flags))
+config.substitutions.append(("%flags", config.test_flags))
+
+config.substitutions.append(("%run", "%t"))

openmp/trunk/libomptarget/deviceRTLs/nvptx/test/lit.site.cfg.in

+@AUTO_GEN_COMMENT@
+
+config.test_c_compiler = "@OPENMP_TEST_C_COMPILER@"
+config.test_cxx_compiler = "@OPENMP_TEST_CXX_COMPILER@"
+config.test_openmp_flags = "@LIBOMPTARGET_NVPTX_TEST_OPENMP_FLAGS@"
+config.test_extra_flags = "@LIBOMPTARGET_NVPTX_TEST_FLAGS@"
+config.binary_dir = "@CMAKE_CURRENT_BINARY_DIR@"
+config.library_dir = "@LIBOMPTARGET_LIBRARY_DIR@"
+config.omp_header_directory = "@LIBOMPTARGET_OPENMP_HEADER_FOLDER@"
+config.omp_host_rtl_directory = "@LIBOMPTARGET_OPENMP_HOST_RTL_FOLDER@"
+config.libomptarget_filecheck = "@OPENMP_FILECHECK_EXECUTABLE@"
+
+# Let the main config do the real work.
+lit_config.load_config(config, "@CMAKE_CURRENT_SOURCE_DIR@/lit.cfg")

openmp/trunk/libomptarget/deviceRTLs/nvptx/test/parallel/thread_limit.c

+// RUN: %compile-run-and-check
+
+#include <stdio.h>
+#include <omp.h>
+
+const int WarpSize = 32;
+const int ThreadLimit = 1 * WarpSize;
+const int NumThreads2 = 2 * WarpSize;
+const int NumThreads3 = 3 * WarpSize;
+const int MaxThreads = 1024;
+
+int main(int argc, char *argv[]) {
+  int check1[MaxThreads];
+  int check2[MaxThreads];
+  int check3[MaxThreads];
+  for (int i = 0; i < MaxThreads; i++) {
+    check1[i] = check2[i] = check3[i] = 0;
+  }
+
+  int threadLimit = -1;
+
+  #pragma omp target teams num_teams(1) thread_limit(ThreadLimit) \
+                           map(check1[:], check2[:], check3[:], threadLimit)
+  {
+    threadLimit = omp_get_thread_limit();
+
+    // All parallel regions should get as many threads as specified by the
+    // thread_limit() clause.
+    #pragma omp parallel
+    {
+      check1[omp_get_thread_num()] += omp_get_num_threads();
+    }
+
+    omp_set_num_threads(NumThreads2);
+    #pragma omp parallel
+    {
+      check2[omp_get_thread_num()] += omp_get_num_threads();
+    }
+
+    #pragma omp parallel num_threads(NumThreads3)
+    {
+      check3[omp_get_thread_num()] += omp_get_num_threads();
+    }
+  }
+
+  // CHECK: threadLimit = 32
+  printf("threadLimit = %d\n", threadLimit);
+
+  // CHECK-NOT: invalid
+  for (int i = 0; i < MaxThreads; i++) {
+    if (i < ThreadLimit) {
+      if (check1[i] != ThreadLimit) {
+        printf("invalid: check1[%d] should be %d, is %d\n", i, ThreadLimit, check1[i]);
+      }
+    } else if (check1[i] != 0) {
+      printf("invalid: check1[%d] should be 0, is %d\n", i, check1[i]);
+    }
+
+    if (i < ThreadLimit) {
+      if (check2[i] != ThreadLimit) {
+        printf("invalid: check2[%d] should be %d, is %d\n", i, ThreadLimit, check2[i]);
+      }
+    } else if (check2[i] != 0) {
+      printf("invalid: check2[%d] should be 0, is %d\n", i, check2[i]);
+    }
+
+    if (i < ThreadLimit) {
+      if (check3[i] != ThreadLimit) {
+        printf("invalid: check3[%d] should be %d, is %d\n", i, ThreadLimit, check3[i]);
+      }
+    } else if (check3[i] != 0) {
+      printf("invalid: check3[%d] should be 0, is %d\n", i, check3[i]);
+    }
+  }
+
+  return 0;
+}

openmp/trunk/libomptarget/test/CMakeLists.txt

 # CMakeLists.txt file for unit testing OpenMP offloading runtime library.
 if(NOT OPENMP_TEST_COMPILER_ID STREQUAL "Clang" OR
    OPENMP_TEST_COMPILER_VERSION VERSION_LESS 6.0.0)
   libomptarget_say("Can only test with Clang compiler in version 6.0.0 or later.")
   libomptarget_warning_say("The check-libomptarget target will not be available!")
   return()
 endif()
 
 if(LIBOMPTARGET_CMAKE_BUILD_TYPE MATCHES debug)
   set(LIBOMPTARGET_DEBUG True)
 else()
   set(LIBOMPTARGET_DEBUG False)
 endif()
 
 add_openmp_testsuite(check-libomptarget "Running libomptarget tests" ${CMAKE_CURRENT_BINARY_DIR} DEPENDS omptarget omp)
 
-if(${OPENMP_STANDALONE_BUILD})
-  set(LIBOMPTARGET_OPENMP_HEADER_FOLDER "${CMAKE_CURRENT_BINARY_DIR}/../../runtime/src" CACHE STRING
-    "Path to folder containing omp.h")
-  set(LIBOMPTARGET_OPENMP_HOST_RTL_FOLDER "${CMAKE_CURRENT_BINARY_DIR}/../../runtime/src" CACHE STRING
-    "Path to folder containing libomp.so")
-else()
-  set(LIBOMPTARGET_OPENMP_HEADER_FOLDER "${LIBOMPTARGET_BINARY_DIR}/../runtime/src")
-endif()
-
 # Configure the lit.site.cfg.in file
 set(AUTO_GEN_COMMENT "## Autogenerated by libomptarget configuration.\n# Do not edit!")
 configure_file(lit.site.cfg.in lit.site.cfg @ONLY)

openmp/trunk/libomptarget/test/lit.site.cfg.in

 config.omp_host_rtl_directory = "@LIBOMPTARGET_OPENMP_HOST_RTL_FOLDER@"
 config.operating_system = "@CMAKE_SYSTEM_NAME@"
 config.libomptarget_all_targets = "@LIBOMPTARGET_ALL_TARGETS@".split()
 config.libomptarget_system_targets = "@LIBOMPTARGET_SYSTEM_TARGETS@".split()
 config.libomptarget_filecheck = "@OPENMP_FILECHECK_EXECUTABLE@"
 config.libomptarget_debug = @LIBOMPTARGET_DEBUG@
 
 # Let the main config do the real work.
-lit_config.load_config(config, "@LIBOMPTARGET_BASE_DIR@/test/lit.cfg")
+lit_config.load_config(config, "@CMAKE_CURRENT_SOURCE_DIR@/lit.cfg")