[libomptarget-nvptx] Add testing infrastructure

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.

Differential Revision: https://reviews.llvm.org/D51687

git-svn-id: https://llvm.org/svn/llvm-project/openmp/trunk@343324 91177308-0d34-0410-b5e6-96231b3b80d8

5 files changed, 191 insertions, 0 deletions

diff --git a/libomptarget/deviceRTLs/nvptx/CMakeLists.txt b/libomptarget/deviceRTLs/nvptx/CMakeLists.txt
index 4fc9ef0..975dee6 100644
--- a/libomptarget/deviceRTLs/nvptx/CMakeLists.txt
+++ b/libomptarget/deviceRTLs/nvptx/CMakeLists.txt
@@ -132,6 +132,9 @@ if(LIBOMPTARGET_DEP_CUDA_FOUND)
       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})
@@ -165,6 +168,7 @@ if(LIBOMPTARGET_DEP_CUDA_FOUND)
       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
@@ -176,6 +180,7 @@ if(LIBOMPTARGET_DEP_CUDA_FOUND)
     endforeach()
   endif()
 
+  add_subdirectory(test)
 else()
   libomptarget_say("Not building CUDA offloading device RTL: CUDA tools not found in the system.")
 endif()
diff --git a/libomptarget/deviceRTLs/nvptx/test/CMakeLists.txt b/libomptarget/deviceRTLs/nvptx/test/CMakeLists.txt
new file mode 100644
index 0000000..33945d1
--- /dev/null
+++ b/libomptarget/deviceRTLs/nvptx/test/CMakeLists.txt
@@ -0,0 +1,26 @@
+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)
diff --git a/libomptarget/deviceRTLs/nvptx/test/lit.cfg b/libomptarget/deviceRTLs/nvptx/test/lit.cfg
new file mode 100644
index 0000000..0774c25
--- /dev/null
+++ b/libomptarget/deviceRTLs/nvptx/test/lit.cfg
@@ -0,0 +1,69 @@
+# -*- 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"))
diff --git a/libomptarget/deviceRTLs/nvptx/test/lit.site.cfg.in b/libomptarget/deviceRTLs/nvptx/test/lit.site.cfg.in
new file mode 100644
index 0000000..d9c14cb
--- /dev/null
+++ b/libomptarget/deviceRTLs/nvptx/test/lit.site.cfg.in
@@ -0,0 +1,14 @@
+@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")
diff --git a/libomptarget/deviceRTLs/nvptx/test/parallel/thread_limit.c b/libomptarget/deviceRTLs/nvptx/test/parallel/thread_limit.c
new file mode 100644
index 0000000..5e40bb5
--- /dev/null
+++ b/libomptarget/deviceRTLs/nvptx/test/parallel/thread_limit.c
@@ -0,0 +1,77 @@
+// 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;
+}