From fea49b2df0172611d0f3624e055979dd31914e21 Mon Sep 17 00:00:00 2001 From: Tobias Ribizel Date: Tue, 17 Nov 2020 21:06:30 +0100 Subject: CTest: Rename CudaMemcheck to CudaSanitizer --- Help/manual/ctest.1.rst | 10 +- Help/variable/CTEST_MEMORYCHECK_TYPE.rst | 2 +- Modules/DartConfiguration.tcl.in | 4 +- Source/CTest/cmCTestMemCheckHandler.cxx | 30 +-- Source/CTest/cmCTestMemCheckHandler.h | 2 +- Tests/RunCMake/CMakeLists.txt | 2 +- .../ctest_memcheck/DummyCudaMemcheck-result.txt | 1 - .../ctest_memcheck/DummyCudaMemcheck-stderr.txt | 1 - .../ctest_memcheck/DummyCudaMemcheck-stdout.txt | 13 - .../ctest_memcheck/DummyCudaSanitizer-result.txt | 1 + .../ctest_memcheck/DummyCudaSanitizer-stderr.txt | 1 + .../ctest_memcheck/DummyCudaSanitizer-stdout.txt | 13 + Tests/RunCMake/ctest_memcheck/RunCMakeTest.cmake | 4 +- .../RunCMake/ctest_memcheck/testCudaMemcheck.cmake | 279 --------------------- .../ctest_memcheck/testCudaSanitizer.cmake | 279 +++++++++++++++++++++ 15 files changed, 321 insertions(+), 321 deletions(-) delete mode 100644 Tests/RunCMake/ctest_memcheck/DummyCudaMemcheck-result.txt delete mode 100644 Tests/RunCMake/ctest_memcheck/DummyCudaMemcheck-stderr.txt delete mode 100644 Tests/RunCMake/ctest_memcheck/DummyCudaMemcheck-stdout.txt create mode 100644 Tests/RunCMake/ctest_memcheck/DummyCudaSanitizer-result.txt create mode 100644 Tests/RunCMake/ctest_memcheck/DummyCudaSanitizer-stderr.txt create mode 100644 Tests/RunCMake/ctest_memcheck/DummyCudaSanitizer-stdout.txt delete mode 100644 Tests/RunCMake/ctest_memcheck/testCudaMemcheck.cmake create mode 100644 Tests/RunCMake/ctest_memcheck/testCudaSanitizer.cmake diff --git a/Help/manual/ctest.1.rst b/Help/manual/ctest.1.rst index e947232..00df24b 100644 --- a/Help/manual/ctest.1.rst +++ b/Help/manual/ctest.1.rst @@ -1142,19 +1142,19 @@ Additional configuration settings include: * `CTest Script`_ variable: none * :module:`CTest` module variable: ``DRMEMORY_COMMAND_OPTIONS`` -``CudaMemcheckCommand`` +``CudaSanitizerCommand`` Specify a ``MemoryCheckCommand`` that is known to be a command-line compatible with cuda-memcheck or compute-sanitizer. * `CTest Script`_ variable: none - * :module:`CTest` module variable: ``CUDA_MEMCHECK_COMMAND`` + * :module:`CTest` module variable: ``CUDA_SANITIZER_COMMAND`` -``CudaMemcheckCommandOptions`` - Specify command-line options to the ``CudaMemcheckCommand`` tool. +``CudaSanitizerCommandOptions`` + Specify command-line options to the ``CudaSanitizerCommand`` tool. They will be placed prior to the test command line. * `CTest Script`_ variable: none - * :module:`CTest` module variable: ``CUDA_MEMCHECK_COMMAND_OPTIONS`` + * :module:`CTest` module variable: ``CUDA_SANITIZER_COMMAND_OPTIONS`` .. _`CTest Submit Step`: diff --git a/Help/variable/CTEST_MEMORYCHECK_TYPE.rst b/Help/variable/CTEST_MEMORYCHECK_TYPE.rst index 5204389..80353a4 100644 --- a/Help/variable/CTEST_MEMORYCHECK_TYPE.rst +++ b/Help/variable/CTEST_MEMORYCHECK_TYPE.rst @@ -6,5 +6,5 @@ CTEST_MEMORYCHECK_TYPE Specify the CTest ``MemoryCheckType`` setting in a :manual:`ctest(1)` dashboard client script. Valid values are ``Valgrind``, ``Purify``, ``BoundsChecker``, ``DrMemory``, -``CudaMemcheck``, ``ThreadSanitizer``, ``AddressSanitizer``, ``LeakSanitizer``, +``CudaSanitizer``, ``ThreadSanitizer``, ``AddressSanitizer``, ``LeakSanitizer``, ``MemorySanitizer`` and ``UndefinedBehaviorSanitizer``. diff --git a/Modules/DartConfiguration.tcl.in b/Modules/DartConfiguration.tcl.in index d318c26..90a56e2 100644 --- a/Modules/DartConfiguration.tcl.in +++ b/Modules/DartConfiguration.tcl.in @@ -71,8 +71,8 @@ ValgrindCommand: @VALGRIND_COMMAND@ ValgrindCommandOptions: @VALGRIND_COMMAND_OPTIONS@ DrMemoryCommand: @DRMEMORY_COMMAND@ DrMemoryCommandOptions: @DRMEMORY_COMMAND_OPTIONS@ -CudaMemcheckCommand: @CUDA_MEMCHECK_COMMAND@ -CudaMemcheckCommandOptions: @CUDA_MEMCHECK_COMMAND_OPTIONS@ +CudaSanitizerCommand: @CUDA_SANITIZER_COMMAND@ +CudaSanitizerCommandOptions: @CUDA_SANITIZER_COMMAND_OPTIONS@ MemoryCheckType: @MEMORYCHECK_TYPE@ MemoryCheckSanitizerOptions: @MEMORYCHECK_SANITIZER_OPTIONS@ MemoryCheckCommand: @MEMORYCHECK_COMMAND@ diff --git a/Source/CTest/cmCTestMemCheckHandler.cxx b/Source/CTest/cmCTestMemCheckHandler.cxx index 73bf764..8a30dc0 100644 --- a/Source/CTest/cmCTestMemCheckHandler.cxx +++ b/Source/CTest/cmCTestMemCheckHandler.cxx @@ -326,8 +326,8 @@ void cmCTestMemCheckHandler::GenerateDartOutput(cmXMLWriter& xml) case cmCTestMemCheckHandler::BOUNDS_CHECKER: xml.Attribute("Checker", "BoundsChecker"); break; - case cmCTestMemCheckHandler::CUDA_MEMCHECK: - xml.Attribute("Checker", "CudaMemcheck"); + case cmCTestMemCheckHandler::CUDA_SANITIZER: + xml.Attribute("Checker", "CudaSanitizer"); break; case cmCTestMemCheckHandler::ADDRESS_SANITIZER: xml.Attribute("Checker", "AddressSanitizer"); @@ -470,7 +470,7 @@ bool cmCTestMemCheckHandler::InitializeMemoryChecking() this->MemoryTesterStyle = cmCTestMemCheckHandler::BOUNDS_CHECKER; } else if (testerName.find("cuda-memcheck") != std::string::npos || testerName.find("compute-sanitizer") != std::string::npos) { - this->MemoryTesterStyle = cmCTestMemCheckHandler::CUDA_MEMCHECK; + this->MemoryTesterStyle = cmCTestMemCheckHandler::CUDA_SANITIZER; } else { this->MemoryTesterStyle = cmCTestMemCheckHandler::UNKNOWN; } @@ -492,10 +492,10 @@ bool cmCTestMemCheckHandler::InitializeMemoryChecking() this->CTest->GetCTestConfiguration("BoundsCheckerCommand"); this->MemoryTesterStyle = cmCTestMemCheckHandler::BOUNDS_CHECKER; } else if (cmSystemTools::FileExists( - this->CTest->GetCTestConfiguration("CudaMemcheckCommand"))) { + this->CTest->GetCTestConfiguration("CudaSanitizerCommand"))) { this->MemoryTester = - this->CTest->GetCTestConfiguration("CudaMemcheckCommand"); - this->MemoryTesterStyle = cmCTestMemCheckHandler::CUDA_MEMCHECK; + this->CTest->GetCTestConfiguration("CudaSanitizerCommand"); + this->MemoryTesterStyle = cmCTestMemCheckHandler::CUDA_SANITIZER; } if (this->CTest->GetCTestConfiguration("MemoryCheckType") == "AddressSanitizer") { @@ -539,8 +539,8 @@ bool cmCTestMemCheckHandler::InitializeMemoryChecking() this->MemoryTesterStyle = cmCTestMemCheckHandler::VALGRIND; } else if (checkType == "DrMemory") { this->MemoryTesterStyle = cmCTestMemCheckHandler::DRMEMORY; - } else if (checkType == "CudaMemcheck") { - this->MemoryTesterStyle = cmCTestMemCheckHandler::CUDA_MEMCHECK; + } else if (checkType == "CudaSanitizer") { + this->MemoryTesterStyle = cmCTestMemCheckHandler::CUDA_SANITIZER; } } if (this->MemoryTester.empty()) { @@ -566,10 +566,10 @@ bool cmCTestMemCheckHandler::InitializeMemoryChecking() .empty()) { memoryTesterOptions = this->CTest->GetCTestConfiguration("DrMemoryCommandOptions"); - } else if (!this->CTest->GetCTestConfiguration("CudaMemcheckCommandOptions") + } else if (!this->CTest->GetCTestConfiguration("CudaSanitizerCommandOptions") .empty()) { memoryTesterOptions = - this->CTest->GetCTestConfiguration("CudaMemcheckCommandOptions"); + this->CTest->GetCTestConfiguration("CudaSanitizerCommandOptions"); } this->MemoryTesterOptions = cmSystemTools::ParseArguments(memoryTesterOptions); @@ -703,8 +703,8 @@ bool cmCTestMemCheckHandler::InitializeMemoryChecking() this->MemoryTesterOptions.emplace_back("/M"); break; } - case cmCTestMemCheckHandler::CUDA_MEMCHECK: { - // cuda-memcheck separates flags from arguments by spaces + case cmCTestMemCheckHandler::CUDA_SANITIZER: { + // cuda sanitizer separates flags from arguments by spaces if (this->MemoryTesterOptions.empty()) { this->MemoryTesterOptions.emplace_back("--tool"); this->MemoryTesterOptions.emplace_back("memcheck"); @@ -800,7 +800,7 @@ bool cmCTestMemCheckHandler::ProcessMemCheckOutput(const std::string& str, return this->ProcessMemCheckSanitizerOutput(str, log, results); case cmCTestMemCheckHandler::BOUNDS_CHECKER: return this->ProcessMemCheckBoundsCheckerOutput(str, log, results); - case cmCTestMemCheckHandler::CUDA_MEMCHECK: + case cmCTestMemCheckHandler::CUDA_SANITIZER: return this->ProcessMemCheckCudaOutput(str, log, results); default: log.append("\nMemory checking style used was: "); @@ -1188,7 +1188,7 @@ bool cmCTestMemCheckHandler::ProcessMemCheckCudaOutput( if (memcheckLine.find(lines[cc])) { cmCTestOptionalLog(this->CTest, DEBUG, - "cuda-memcheck line " << lines[cc] << std::endl, + "cuda sanitizer line " << lines[cc] << std::endl, this->Quiet); int failure = -1; auto& line = lines[cc]; @@ -1219,7 +1219,7 @@ bool cmCTestMemCheckHandler::ProcessMemCheckCudaOutput( nonMemcheckOutput.push_back(cc); } } - // Now put all all the non cuda-memcheck output into the test output + // Now put all all the non cuda sanitizer output into the test output // This should be last in case it gets truncated by the output // limiting code for (std::string::size_type i : nonMemcheckOutput) { diff --git a/Source/CTest/cmCTestMemCheckHandler.h b/Source/CTest/cmCTestMemCheckHandler.h index 6ef5d20..7ab00db 100644 --- a/Source/CTest/cmCTestMemCheckHandler.h +++ b/Source/CTest/cmCTestMemCheckHandler.h @@ -45,7 +45,7 @@ private: DRMEMORY, BOUNDS_CHECKER, // checkers after here do not use the standard error list - CUDA_MEMCHECK, + CUDA_SANITIZER, ADDRESS_SANITIZER, LEAK_SANITIZER, THREAD_SANITIZER, diff --git a/Tests/RunCMake/CMakeLists.txt b/Tests/RunCMake/CMakeLists.txt index 06dd859..7000b79 100644 --- a/Tests/RunCMake/CMakeLists.txt +++ b/Tests/RunCMake/CMakeLists.txt @@ -167,7 +167,7 @@ if(NOT CMake_TEST_EXTERNAL_CMAKE) -DPSEUDO_BC=$ -DPSEUDO_PURIFY=$ -DPSEUDO_VALGRIND=$ - -DPSEUDO_CUDA_MEMCHECK=$ + -DPSEUDO_CUDA_SANITIZER=$ -DPSEUDO_BC_NOLOG=$ -DPSEUDO_PURIFY_NOLOG=$ -DPSEUDO_VALGRIND_NOLOG=$ diff --git a/Tests/RunCMake/ctest_memcheck/DummyCudaMemcheck-result.txt b/Tests/RunCMake/ctest_memcheck/DummyCudaMemcheck-result.txt deleted file mode 100644 index 573541a..0000000 --- a/Tests/RunCMake/ctest_memcheck/DummyCudaMemcheck-result.txt +++ /dev/null @@ -1 +0,0 @@ -0 diff --git a/Tests/RunCMake/ctest_memcheck/DummyCudaMemcheck-stderr.txt b/Tests/RunCMake/ctest_memcheck/DummyCudaMemcheck-stderr.txt deleted file mode 100644 index d302b5c..0000000 --- a/Tests/RunCMake/ctest_memcheck/DummyCudaMemcheck-stderr.txt +++ /dev/null @@ -1 +0,0 @@ -Defect count: 23 diff --git a/Tests/RunCMake/ctest_memcheck/DummyCudaMemcheck-stdout.txt b/Tests/RunCMake/ctest_memcheck/DummyCudaMemcheck-stdout.txt deleted file mode 100644 index 034ee1e..0000000 --- a/Tests/RunCMake/ctest_memcheck/DummyCudaMemcheck-stdout.txt +++ /dev/null @@ -1,13 +0,0 @@ -Memory checking results: -Uninitialized __global__ memory read - 1 -Unused memory - 1 -Host API memory access error - 1 -Barrier error - 2 -Invalid __global__ read - 1 -cudaErrorLaunchFailure - 2 -Fatal UVM GPU fault - 1 -Fatal UVM CPU fault - 1 -Memory leak - 1 -Potential WAR hazard detected - 4 -Potential RAW hazard detected - 4 -Race reported - 4 diff --git a/Tests/RunCMake/ctest_memcheck/DummyCudaSanitizer-result.txt b/Tests/RunCMake/ctest_memcheck/DummyCudaSanitizer-result.txt new file mode 100644 index 0000000..573541a --- /dev/null +++ b/Tests/RunCMake/ctest_memcheck/DummyCudaSanitizer-result.txt @@ -0,0 +1 @@ +0 diff --git a/Tests/RunCMake/ctest_memcheck/DummyCudaSanitizer-stderr.txt b/Tests/RunCMake/ctest_memcheck/DummyCudaSanitizer-stderr.txt new file mode 100644 index 0000000..d302b5c --- /dev/null +++ b/Tests/RunCMake/ctest_memcheck/DummyCudaSanitizer-stderr.txt @@ -0,0 +1 @@ +Defect count: 23 diff --git a/Tests/RunCMake/ctest_memcheck/DummyCudaSanitizer-stdout.txt b/Tests/RunCMake/ctest_memcheck/DummyCudaSanitizer-stdout.txt new file mode 100644 index 0000000..034ee1e --- /dev/null +++ b/Tests/RunCMake/ctest_memcheck/DummyCudaSanitizer-stdout.txt @@ -0,0 +1,13 @@ +Memory checking results: +Uninitialized __global__ memory read - 1 +Unused memory - 1 +Host API memory access error - 1 +Barrier error - 2 +Invalid __global__ read - 1 +cudaErrorLaunchFailure - 2 +Fatal UVM GPU fault - 1 +Fatal UVM CPU fault - 1 +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 2b3165b..6e0a91c 100644 --- a/Tests/RunCMake/ctest_memcheck/RunCMakeTest.cmake +++ b/Tests/RunCMake/ctest_memcheck/RunCMakeTest.cmake @@ -179,11 +179,11 @@ unset(CTEST_SUFFIX_CODE) #----------------------------------------------------------------------------- set(CMAKELISTS_EXTRA_CODE "add_test(NAME TestSan COMMAND \"${CMAKE_COMMAND}\" --P \"${RunCMake_SOURCE_DIR}/testCudaMemcheck.cmake\") +-P \"${RunCMake_SOURCE_DIR}/testCudaSanitizer.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}") +run_mc_test(DummyCudaSanitizer "${PSEUDO_CUDA_SANITIZER}") 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 deleted file mode 100644 index adc7a1a..0000000 --- a/Tests/RunCMake/ctest_memcheck/testCudaMemcheck.cmake +++ /dev/null @@ -1,279 +0,0 @@ -# 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] -========= -========= Host API memory access error at host access to 0x1303fd1400 of size 25600 bytes -========= Uninitialized access at 0x1303fd4600 on access by cudaMemcopy source. -========= Saved host backtrace up to driver entry point at error -========= Host Frame:/usr/lib/x86_64-linux-gnu/libcuda.so.1 (cuMemcpyDtoH_v2 + 0x1ec) [0x29200c] -========= Host Frame:/usr/local/cuda/targets/x86_64-linux/lib/libcudart.so.10.1 [0x38aaa] -========= Host Frame:/usr/local/cuda/targets/x86_64-linux/lib/libcudart.so.10.1 [0x18946] -========= Host Frame:/usr/local/cuda/targets/x86_64-linux/lib/libcudart.so.10.1 (cudaMemcpy + 0x1a2) [0x3b8c2] -========= Host Frame:/something/somewhere [0xcafe] -========= -========= 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] -========= -========= Fatal UVM GPU fault of type invalid pde due to invalid address -========= during atomic access to address 0x20be00000 -========= -========= Fatal UVM CPU fault due to invalid operation -========= during read access to address 0x1357c92000 -========= -========= 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) -") diff --git a/Tests/RunCMake/ctest_memcheck/testCudaSanitizer.cmake b/Tests/RunCMake/ctest_memcheck/testCudaSanitizer.cmake new file mode 100644 index 0000000..adc7a1a --- /dev/null +++ b/Tests/RunCMake/ctest_memcheck/testCudaSanitizer.cmake @@ -0,0 +1,279 @@ +# 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] +========= +========= Host API memory access error at host access to 0x1303fd1400 of size 25600 bytes +========= Uninitialized access at 0x1303fd4600 on access by cudaMemcopy source. +========= Saved host backtrace up to driver entry point at error +========= Host Frame:/usr/lib/x86_64-linux-gnu/libcuda.so.1 (cuMemcpyDtoH_v2 + 0x1ec) [0x29200c] +========= Host Frame:/usr/local/cuda/targets/x86_64-linux/lib/libcudart.so.10.1 [0x38aaa] +========= Host Frame:/usr/local/cuda/targets/x86_64-linux/lib/libcudart.so.10.1 [0x18946] +========= Host Frame:/usr/local/cuda/targets/x86_64-linux/lib/libcudart.so.10.1 (cudaMemcpy + 0x1a2) [0x3b8c2] +========= Host Frame:/something/somewhere [0xcafe] +========= +========= 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] +========= +========= Fatal UVM GPU fault of type invalid pde due to invalid address +========= during atomic access to address 0x20be00000 +========= +========= Fatal UVM CPU fault due to invalid operation +========= during read access to address 0x1357c92000 +========= +========= 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) +") -- cgit v0.12