Merge topic 'cuda-memcheck'

cee92a9fb0 Help: add release notes for CTest cuda-memcheck support
f38e4a1871 Tests: Add cases for CTest cuda-memcheck parser
fe062800f0 CTest: add cuda-memcheck support

Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: tcojean <terry.cojean@kit.edu>
Merge-request: !4952

8 files changed, 324 insertions, 2 deletions

diff --git a/Tests/CMakeLib/PseudoMemcheck/CMakeLists.txt b/Tests/CMakeLib/PseudoMemcheck/CMakeLists.txt
index 7c84ee1..4bef6c5 100644
--- a/Tests/CMakeLib/PseudoMemcheck/CMakeLists.txt
+++ b/Tests/CMakeLib/PseudoMemcheck/CMakeLists.txt
@@ -15,6 +15,9 @@ target_link_libraries(pseudo_purify CMakeLib)
 add_executable(pseudo_BC "${CMAKE_CURRENT_BINARY_DIR}/ret0.cxx")
 set_target_properties(pseudo_BC PROPERTIES OUTPUT_NAME BC)
 target_link_libraries(pseudo_BC CMakeLib)
+add_executable(pseudo_cuda-memcheck "${CMAKE_CURRENT_BINARY_DIR}/ret0.cxx")
+set_target_properties(pseudo_cuda-memcheck PROPERTIES OUTPUT_NAME cuda-memcheck)
+target_link_libraries(pseudo_cuda-memcheck CMakeLib)
 
 # binary to be used as pre- and post-memcheck command that fails
 add_executable(memcheck_fail "${CMAKE_CURRENT_BINARY_DIR}/ret1.cxx")
diff --git a/Tests/CMakeLib/PseudoMemcheck/memtester.cxx.in b/Tests/CMakeLib/PseudoMemcheck/memtester.cxx.in
index 3183bc0..f37ad59 100644
--- a/Tests/CMakeLib/PseudoMemcheck/memtester.cxx.in
+++ b/Tests/CMakeLib/PseudoMemcheck/memtester.cxx.in
@@ -1,8 +1,14 @@
-#include <cmSystemTools.h>
-#include "cmsys/Encoding.hxx"
 #include <string>
+#include <vector>
+
+#include "cmsys/Encoding.hxx"
+
+#include <cmSystemTools.h>
 
+// clang-format off
 #define RETVAL @_retval@
+#define CMAKE_COMMAND "@CMAKE_COMMAND@"
+// clang-format on
 
 int main(int ac, char** av)
 {
@@ -14,6 +20,9 @@ int main(int ac, char** av)
   std::string exename = argv[0];
   std::string logarg;
   bool nextarg = false;
+  // execute the part after the last argument?
+  // the logfile path gets passed as environment variable PSEUDO_LOGFILE
+  bool exec = false;
 
   if (exename.find("valgrind") != std::string::npos) {
     logarg = "--log-file=";
@@ -26,6 +35,10 @@ int main(int ac, char** av)
   } else if (exename.find("BC") != std::string::npos) {
     nextarg = true;
     logarg = "/X";
+  } else if (exename.find("cuda-memcheck") != std::string::npos) {
+    nextarg = true;
+    exec = true;
+    logarg = "--log-file";
   }
 
   if (!logarg.empty()) {
@@ -45,8 +58,25 @@ int main(int ac, char** av)
       }
     }
 
+    // find the last argument position
+    int lastarg_pos = 1;
+    for (int i = 1; i < argc; ++i) {
+      std::string arg = argv[i];
+      if (arg.find("--") == 0) {
+        lastarg_pos = i;
+      }
+    }
+
     if (!logfile.empty()) {
       cmSystemTools::Touch(logfile, true);
+      // execute everything after the last argument with additional environment
+      int callarg_pos = lastarg_pos + (nextarg ? 2 : 1);
+      if (exec && callarg_pos < argc) {
+        std::vector<std::string> callargs{ CMAKE_COMMAND, "-E", "env",
+                                           "PSEUDO_LOGFILE=" + logfile };
+        callargs.insert(callargs.end(), &argv[callarg_pos], &argv[argc]);
+        cmSystemTools::RunSingleCommand(callargs);
+      }
     }
   }
 
diff --git a/Tests/RunCMake/CMakeLists.txt b/Tests/RunCMake/CMakeLists.txt
index 871e18f..b037a6d 100644
--- a/Tests/RunCMake/CMakeLists.txt
+++ b/Tests/RunCMake/CMakeLists.txt
@@ -163,6 +163,7 @@ if(NOT CMake_TEST_EXTERNAL_CMAKE)
     -DPSEUDO_BC=$<TARGET_FILE:pseudo_BC>
     -DPSEUDO_PURIFY=$<TARGET_FILE:pseudo_purify>
     -DPSEUDO_VALGRIND=$<TARGET_FILE:pseudo_valgrind>
+    -DPSEUDO_CUDA_MEMCHECK=$<TARGET_FILE:pseudo_cuda-memcheck>
     -DPSEUDO_BC_NOLOG=$<TARGET_FILE:pseudonl_BC>
     -DPSEUDO_PURIFY_NOLOG=$<TARGET_FILE:pseudonl_purify>
     -DPSEUDO_VALGRIND_NOLOG=$<TARGET_FILE:pseudonl_valgrind>
diff --git a/Tests/RunCMake/ctest_memcheck/DummyCudaMemcheck-result.txt b/Tests/RunCMake/ctest_memcheck/DummyCudaMemcheck-result.txt
new file mode 100644
index 0000000..573541a
--- /dev/null
+++ b/Tests/RunCMake/ctest_memcheck/DummyCudaMemcheck-result.txt
@@ -0,0 +1 @@
+0
diff --git a/Tests/RunCMake/ctest_memcheck/DummyCudaMemcheck-stderr.txt b/Tests/RunCMake/ctest_memcheck/DummyCudaMemcheck-stderr.txt
new file mode 100644
index 0000000..31dedd2
--- /dev/null
+++ b/Tests/RunCMake/ctest_memcheck/DummyCudaMemcheck-stderr.txt
@@ -0,0 +1 @@
+Defect count: 20
diff --git a/Tests/RunCMake/ctest_memcheck/DummyCudaMemcheck-stdout.txt b/Tests/RunCMake/ctest_memcheck/DummyCudaMemcheck-stdout.txt
new file mode 100644
index 0000000..aa3e698
--- /dev/null
+++ b/Tests/RunCMake/ctest_memcheck/DummyCudaMemcheck-stdout.txt
@@ -0,0 +1,10 @@
+Memory checking results:
+Uninitialized __global__ memory read - 1
+Unused memory - 1
+Barrier error - 2
+Invalid __global__ read - 1
+cudaErrorLaunchFailure - 2
+Memory leak - 1
+Potential WAR hazard detected - 4
+Potential RAW hazard detected - 4
+Race reported - 4
diff --git a/Tests/RunCMake/ctest_memcheck/RunCMakeTest.cmake b/Tests/RunCMake/ctest_memcheck/RunCMakeTest.cmake
index ab4c5ab..2b3165b 100644
--- a/Tests/RunCMake/ctest_memcheck/RunCMakeTest.cmake
+++ b/Tests/RunCMake/ctest_memcheck/RunCMakeTest.cmake
@@ -175,3 +175,15 @@ unset(CMAKELISTS_EXTRA_CODE)
 unset(CTEST_EXTRA_CODE)
 unset(CTEST_MEMCHECK_ARGS)
 unset(CTEST_SUFFIX_CODE)
+
+#-----------------------------------------------------------------------------
+set(CMAKELISTS_EXTRA_CODE
+"add_test(NAME TestSan COMMAND \"${CMAKE_COMMAND}\"
+-P \"${RunCMake_SOURCE_DIR}/testCudaMemcheck.cmake\")
+")
+set(CTEST_SUFFIX_CODE "message(\"Defect count: \${defect_count}\")")
+set(CTEST_MEMCHECK_ARGS "DEFECT_COUNT defect_count")
+run_mc_test(DummyCudaMemcheck "${PSEUDO_CUDA_MEMCHECK}")
+unset(CTEST_MEMCHECK_ARGS)
+unset(CTEST_SUFFIX_CODE)
+unset(CTEST_EXTRA_CODE)
diff --git a/Tests/RunCMake/ctest_memcheck/testCudaMemcheck.cmake b/Tests/RunCMake/ctest_memcheck/testCudaMemcheck.cmake
new file mode 100644
index 0000000..d7d8213
--- /dev/null
+++ b/Tests/RunCMake/ctest_memcheck/testCudaMemcheck.cmake
@@ -0,0 +1,264 @@
+# this file simulates an execution of cuda-memcheck
+
+set(LOG_FILE "$ENV{PSEUDO_LOGFILE}")
+message("LOG_FILE=[${LOG_FILE}]")
+
+# clear the log file
+file(REMOVE "${LOG_FILE}")
+
+# create an error of each type of sanitizer tool and failure
+
+# initcheck
+file(APPEND "${LOG_FILE}"
+"========= CUDA-MEMCHECK
+========= Uninitialized __global__ memory read of size 4
+=========     at 0x00000020 in test(int*, int*)
+=========     by thread (0,0,0) in block (0,0,0)
+=========     Address 0x1303d80000
+=========     Saved host backtrace up to driver entry point
+=========     Host Frame:/lib64/libcuda.so.1 (cuLaunchKernel + 0x346) [0x297db6]
+=========     Host Frame:./uninit-read [0x101d9]
+=========     Host Frame:./uninit-read [0x10267]
+=========     Host Frame:./uninit-read [0x465b5]
+=========     Host Frame:./uninit-read [0x3342]
+=========     Host Frame:./uninit-read [0x3143]
+=========     Host Frame:/lib64/libc.so.6 (__libc_start_main + 0xf5) [0x22505]
+=========     Host Frame:./uninit-read [0x31e2]
+=========
+========= Unused memory in allocation 0x1303d80000 of size 16 bytes
+=========     Not written any memory.
+=========     100.00% of allocation were unused.
+=========     Saved host backtrace up to driver entry point
+=========     Host Frame:/lib64/libcuda.so.1 (cuMemAlloc_v2 + 0x1b7) [0x26ec97]
+=========     Host Frame:./uninit-read [0x2bbd3]
+=========     Host Frame:./uninit-read [0x71ab]
+=========     Host Frame:./uninit-read [0x3c84f]
+=========     Host Frame:./uninit-read [0x3111]
+=========     Host Frame:/lib64/libc.so.6 (__libc_start_main + 0xf5) [0x22505]
+=========     Host Frame:./uninit-read [0x31e2]
+=========
+========= ERROR SUMMARY: 2 errors
+")
+
+
+# synccheck
+file(APPEND "${LOG_FILE}"
+"========= CUDA-MEMCHECK
+========= Barrier error detected. Divergent thread(s) in warp
+=========     at 0x00000058 in test(int*, int*)
+=========     by thread (1,0,0) in block (0,0,0)
+=========     Device Frame:test(int*, int*) (test(int*, int*) : 0x60)
+=========     Saved host backtrace up to driver entry point at kernel launch time
+=========     Host Frame:/lib64/libcuda.so.1 (cuLaunchKernel + 0x346) [0x297db6]
+=========     Host Frame:./sync [0x101d9]
+=========     Host Frame:./sync [0x10267]
+=========     Host Frame:./sync [0x465b5]
+=========     Host Frame:./sync [0x3342]
+=========     Host Frame:./sync [0x314a]
+=========     Host Frame:/lib64/libc.so.6 (__libc_start_main + 0xf5) [0x22505]
+=========     Host Frame:./sync [0x31e2]
+=========
+========= Barrier error detected. Divergent thread(s) in warp
+=========     at 0x00000058 in test(int*, int*)
+=========     by thread (0,0,0) in block (0,0,0)
+=========     Device Frame:test(int*, int*) (test(int*, int*) : 0x60)
+=========     Saved host backtrace up to driver entry point at kernel launch time
+=========     Host Frame:/lib64/libcuda.so.1 (cuLaunchKernel + 0x346) [0x297db6]
+=========     Host Frame:./sync [0x101d9]
+=========     Host Frame:./sync [0x10267]
+=========     Host Frame:./sync [0x465b5]
+=========     Host Frame:./sync [0x3342]
+=========     Host Frame:./sync [0x314a]
+=========     Host Frame:/lib64/libc.so.6 (__libc_start_main + 0xf5) [0x22505]
+=========     Host Frame:./sync [0x31e2]
+=========
+========= ERROR SUMMARY: 2 errors
+")
+
+# memcheck
+file(APPEND "${LOG_FILE}"
+"========= CUDA-MEMCHECK
+========= Invalid __global__ read of size 4
+=========     at 0x00000020 in test(int*, int*)
+=========     by thread (0,0,0) in block (0,0,0)
+=========     Address 0x00000000 is out of bounds
+=========     Saved host backtrace up to driver entry point at kernel launch time
+=========     Host Frame:/lib64/libcuda.so.1 (cuLaunchKernel + 0x346) [0x297db6]
+=========     Host Frame:./invalid-read [0x101d9]
+=========     Host Frame:./invalid-read [0x10267]
+=========     Host Frame:./invalid-read [0x465b5]
+=========     Host Frame:./invalid-read [0x3342]
+=========     Host Frame:./invalid-read [0x3142]
+=========     Host Frame:/lib64/libc.so.6 (__libc_start_main + 0xf5) [0x22505]
+=========     Host Frame:./invalid-read [0x31e2]
+=========
+========= Program hit cudaErrorLaunchFailure (error 719) due to \"unspecified launch failure\" on CUDA API call to cudaDeviceSynchronize.
+=========     Saved host backtrace up to driver entry point at error
+=========     Host Frame:/lib64/libcuda.so.1 [0x3ac5a3]
+=========     Host Frame:./invalid-read [0x2e576]
+=========     Host Frame:./invalid-read [0x3147]
+=========     Host Frame:/lib64/libc.so.6 (__libc_start_main + 0xf5) [0x22505]
+=========     Host Frame:./invalid-read [0x31e2]
+=========
+========= Program hit cudaErrorLaunchFailure (error 719) due to \"unspecified launch failure\" on CUDA API call to cudaFree.
+=========     Saved host backtrace up to driver entry point at error
+=========     Host Frame:/lib64/libcuda.so.1 [0x3ac5a3]
+=========     Host Frame:./invalid-read [0x3c106]
+=========     Host Frame:./invalid-read [0x3150]
+=========     Host Frame:/lib64/libc.so.6 (__libc_start_main + 0xf5) [0x22505]
+=========     Host Frame:./invalid-read [0x31e2]
+=========
+========= LEAK SUMMARY: 0 bytes leaked in 0 allocations
+========= ERROR SUMMARY: 3 errors
+")
+
+# memcheck with leak-check full
+file(APPEND "${LOG_FILE}"
+"========= CUDA-MEMCHECK
+========= Leaked 10 bytes at 0x1303d80000
+=========     Saved host backtrace up to driver entry point at cudaMalloc time
+=========     Host Frame:/lib64/libcuda.so.1 (cuMemAlloc_v2 + 0x1b7) [0x26ec97]
+=========     Host Frame:./leak [0x2bab3]
+=========     Host Frame:./leak [0x708b]
+=========     Host Frame:./leak [0x3c72f]
+=========     Host Frame:./leak [0x3113]
+=========     Host Frame:/lib64/libc.so.6 (__libc_start_main + 0xf5) [0x22505]
+=========     Host Frame:./leak [0x3174]
+=========
+========= LEAK SUMMARY: 10 bytes leaked in 1 allocations
+========= ERROR SUMMARY: 1 error
+")
+
+# racecheck with racecheck-report all
+file(APPEND "${LOG_FILE}"
+"========= CUDA-MEMCHECK
+========= WARN:(Warp Level Programming) Potential WAR hazard detected at __shared__ 0x3 in block (0, 0, 0) :
+=========     Read Thread (31, 0, 0) at 0x00000170 in ./race.cu:4:test(int*, int*)
+=========     Write Thread (0, 0, 0) at 0x000001a8 in ./race.cu:4:test(int*, int*)
+=========     Current Value : 0, Incoming Value : 0
+=========     Saved host backtrace up to driver entry point at kernel launch time
+=========     Host Frame:/lib64/libcuda.so.1 (cuLaunchKernel + 0x346) [0x297db6]
+=========     Host Frame:./race [0x101d9]
+=========     Host Frame:./race [0x10267]
+=========     Host Frame:./race [0x465b5]
+=========     Host Frame:./race [0x3342]
+=========     Host Frame:./race [0x314a]
+=========     Host Frame:/lib64/libc.so.6 (__libc_start_main + 0xf5) [0x22505]
+=========     Host Frame:./race [0x31e2]
+=========
+========= WARN:(Warp Level Programming) Potential WAR hazard detected at __shared__ 0x2 in block (0, 0, 0) :
+=========     Read Thread (31, 0, 0) at 0x00000170 in ./race.cu:4:test(int*, int*)
+=========     Write Thread (0, 0, 0) at 0x000001a8 in ./race.cu:4:test(int*, int*)
+=========     Current Value : 0, Incoming Value : 0
+=========     Saved host backtrace up to driver entry point at kernel launch time
+=========     Host Frame:/lib64/libcuda.so.1 (cuLaunchKernel + 0x346) [0x297db6]
+=========     Host Frame:./race [0x101d9]
+=========     Host Frame:./race [0x10267]
+=========     Host Frame:./race [0x465b5]
+=========     Host Frame:./race [0x3342]
+=========     Host Frame:./race [0x314a]
+=========     Host Frame:/lib64/libc.so.6 (__libc_start_main + 0xf5) [0x22505]
+=========     Host Frame:./race [0x31e2]
+=========
+========= WARN:(Warp Level Programming) Potential WAR hazard detected at __shared__ 0x1 in block (0, 0, 0) :
+=========     Read Thread (31, 0, 0) at 0x00000170 in ./race.cu:4:test(int*, int*)
+=========     Write Thread (0, 0, 0) at 0x000001a8 in ./race.cu:4:test(int*, int*)
+=========     Current Value : 0, Incoming Value : 0
+=========     Saved host backtrace up to driver entry point at kernel launch time
+=========     Host Frame:/lib64/libcuda.so.1 (cuLaunchKernel + 0x346) [0x297db6]
+=========     Host Frame:./race [0x101d9]
+=========     Host Frame:./race [0x10267]
+=========     Host Frame:./race [0x465b5]
+=========     Host Frame:./race [0x3342]
+=========     Host Frame:./race [0x314a]
+=========     Host Frame:/lib64/libc.so.6 (__libc_start_main + 0xf5) [0x22505]
+=========     Host Frame:./race [0x31e2]
+=========
+========= WARN:(Warp Level Programming) Potential WAR hazard detected at __shared__ 0x0 in block (0, 0, 0) :
+=========     Read Thread (31, 0, 0) at 0x00000170 in ./race.cu:4:test(int*, int*)
+=========     Write Thread (0, 0, 0) at 0x000001a8 in ./race.cu:4:test(int*, int*)
+=========     Current Value : 0, Incoming Value : 1
+=========     Saved host backtrace up to driver entry point at kernel launch time
+=========     Host Frame:/lib64/libcuda.so.1 (cuLaunchKernel + 0x346) [0x297db6]
+=========     Host Frame:./race [0x101d9]
+=========     Host Frame:./race [0x10267]
+=========     Host Frame:./race [0x465b5]
+=========     Host Frame:./race [0x3342]
+=========     Host Frame:./race [0x314a]
+=========     Host Frame:/lib64/libc.so.6 (__libc_start_main + 0xf5) [0x22505]
+=========     Host Frame:./race [0x31e2]
+=========
+========= WARN:(Warp Level Programming) Potential RAW hazard detected at __shared__ 0x3 in block (0, 0, 0) :
+=========     Write Thread (31, 0, 0) at 0x00000148 in ./race.cu:3:test(int*, int*)
+=========     Read Thread (0, 0, 0) at 0x00000170 in ./race.cu:4:test(int*, int*)
+=========     Current Value : 0
+=========     Saved host backtrace up to driver entry point at kernel launch time
+=========     Host Frame:/lib64/libcuda.so.1 (cuLaunchKernel + 0x346) [0x297db6]
+=========     Host Frame:./race [0x101d9]
+=========     Host Frame:./race [0x10267]
+=========     Host Frame:./race [0x465b5]
+=========     Host Frame:./race [0x3342]
+=========     Host Frame:./race [0x314a]
+=========     Host Frame:/lib64/libc.so.6 (__libc_start_main + 0xf5) [0x22505]
+=========     Host Frame:./race [0x31e2]
+=========
+========= WARN:(Warp Level Programming) Potential RAW hazard detected at __shared__ 0x2 in block (0, 0, 0) :
+=========     Write Thread (31, 0, 0) at 0x00000148 in ./race.cu:3:test(int*, int*)
+=========     Read Thread (0, 0, 0) at 0x00000170 in ./race.cu:4:test(int*, int*)
+=========     Current Value : 0
+=========     Saved host backtrace up to driver entry point at kernel launch time
+=========     Host Frame:/lib64/libcuda.so.1 (cuLaunchKernel + 0x346) [0x297db6]
+=========     Host Frame:./race [0x101d9]
+=========     Host Frame:./race [0x10267]
+=========     Host Frame:./race [0x465b5]
+=========     Host Frame:./race [0x3342]
+=========     Host Frame:./race [0x314a]
+=========     Host Frame:/lib64/libc.so.6 (__libc_start_main + 0xf5) [0x22505]
+=========     Host Frame:./race [0x31e2]
+=========
+========= WARN:(Warp Level Programming) Potential RAW hazard detected at __shared__ 0x1 in block (0, 0, 0) :
+=========     Write Thread (31, 0, 0) at 0x00000148 in ./race.cu:3:test(int*, int*)
+=========     Read Thread (0, 0, 0) at 0x00000170 in ./race.cu:4:test(int*, int*)
+=========     Current Value : 0
+=========     Saved host backtrace up to driver entry point at kernel launch time
+=========     Host Frame:/lib64/libcuda.so.1 (cuLaunchKernel + 0x346) [0x297db6]
+=========     Host Frame:./race [0x101d9]
+=========     Host Frame:./race [0x10267]
+=========     Host Frame:./race [0x465b5]
+=========     Host Frame:./race [0x3342]
+=========     Host Frame:./race [0x314a]
+=========     Host Frame:/lib64/libc.so.6 (__libc_start_main + 0xf5) [0x22505]
+=========     Host Frame:./race [0x31e2]
+=========
+========= WARN:(Warp Level Programming) Potential RAW hazard detected at __shared__ 0x0 in block (0, 0, 0) :
+=========     Write Thread (31, 0, 0) at 0x00000148 in ./race.cu:3:test(int*, int*)
+=========     Read Thread (0, 0, 0) at 0x00000170 in ./race.cu:4:test(int*, int*)
+=========     Current Value : 0
+=========     Saved host backtrace up to driver entry point at kernel launch time
+=========     Host Frame:/lib64/libcuda.so.1 (cuLaunchKernel + 0x346) [0x297db6]
+=========     Host Frame:./race [0x101d9]
+=========     Host Frame:./race [0x10267]
+=========     Host Frame:./race [0x465b5]
+=========     Host Frame:./race [0x3342]
+=========     Host Frame:./race [0x314a]
+=========     Host Frame:/lib64/libc.so.6 (__libc_start_main + 0xf5) [0x22505]
+=========     Host Frame:./race [0x31e2]
+=========
+========= WARN: Race reported between Read access at 0x00000170 in ./race.cu:4:test(int*, int*)
+=========     and Write access at 0x00000148 in ./race.cu:3:test(int*, int*) [4 hazards]
+=========     and Write access at 0x000001a8 in ./race.cu:4:test(int*, int*) [4 hazards]
+=========
+========= WARN: Race reported between Write access at 0x00000148 in ./race.cu:3:test(int*, int*)
+=========     and Write access at 0x00000148 in ./race.cu:3:test(int*, int*) [124 hazards]
+=========     and Read access at 0x00000170 in ./race.cu:4:test(int*, int*) [4 hazards]
+=========
+========= WARN: Race reported between Write access at 0x000001a8 in ./race.cu:4:test(int*, int*)
+=========     and Write access at 0x000001a8 in ./race.cu:4:test(int*, int*) [124 hazards]
+=========     and Read access at 0x00000170 in ./race.cu:4:test(int*, int*) [4 hazards]
+=========
+========= WARN: Race reported between Write access at 0x00000148 in ./race.cu:3:test(int*, int*)
+=========     and Write access at 0x00000148 in ./race.cu:3:test(int*, int*) [124 hazards]
+=========     and Read access at 0x00000170 in ./race.cu:4:test(int*, int*) [4 hazards]
+=========
+========= RACECHECK SUMMARY: 12 hazards displayed (0 errors, 12 warnings)
+")