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

6 files changed, 289 insertions, 0 deletions

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)
+")