Sanitizer-lite: memcheck / initcheck / racecheck / synccheck#

A CPU-single-thread simulator sees the whole interleaving, so it can flag bugs a real GPU hides — approximating NVIDIA compute-sanitizer with no GPU. All of it is behind SIM_CHECK (define it before including this header; only the CPU build does). Access the memory you want checked through SIM_LD / SIM_ST; the query API reports how many faults each checker saw. Without SIM_CHECK every hook is a no-op, SIM_LD/ST are a plain p[i], and the counters read zero.

Defined in code/simulator.h.

This page documents 17 symbols: 2 macros, 15 functions.

Macros#

Signature

Description

Availability

Location

Example

SIM_LD(p, i)

Sanitizer-lite hooks are a CPU-sim-only feature (real GPUs use NVIDIA compute-sanitizer).

sim + nvcc

simulator.h:3767

View

SIM_ST(p, i, v)

Expands to (::simdetail::checkedStore((p), (i), (v)))

sim + nvcc

simulator.h:3768

View

Functions#

Signature

Description

Availability

Location

Example

inline T checkedLoad(const T* p, size_t i)

sim only

simulator.h:3721

inline void checkedStore(T* p, size_t i, T v)

sim only

simulator.h:3729

inline int& initErr()

sim only

simulator.h:3693

inline std::unordered_set<uintptr_t>& initSet()

sim only

simulator.h:3690

inline bool inTrackedRange(const void* p, size_t n)

Does [p, p+n) lie inside one tracked allocation (or the dynamic-shared buffer)?

sim only

simulator.h:3698

inline int& memErr()

sim only

simulator.h:3692

inline int& raceErr()

sim only

simulator.h:3694

inline std::unordered_map<uintptr_t, AccessRec>& raceMap()

sim only

simulator.h:3689

inline void raceObserve(uintptr_t a, bool isWrite)

Record an access at byte address a this epoch; flag a data race if another thread already touched it this epoch and at least one access is a write.

sim only

simulator.h:3713

inline void simCheckReset()

sim + nvcc

simulator.h:3763

View

inline int simInitcheckErrors()

sim + nvcc

simulator.h:3760

View

inline int simMemcheckErrors()

Public query API: number of faults each checker has seen since the last simCheckReset().

sim + nvcc

simulator.h:3759

View

inline int simRacecheckErrors()

sim + nvcc

simulator.h:3761

View

inline int simSynccheckErrors()

sim + nvcc

simulator.h:3762

View

inline int& syncErr()

sim only

simulator.h:3695

Examples#

Code from the repository’s example corpus (code\*.cu); the highlighted line is the call site. These blocks are what the View links above open.

SIM_LD (test_sanitizer.cu:38)
28        }                                                                    \
29    } while (0)
30
31constexpr int BLK = 32;
32
33// memcheck: out-of-bounds store vs. in-bounds store.
34__global__ void kOOB(int* out, int n)      { SIM_ST(out, threadIdx.x + n, 1); }
35__global__ void kInBounds(int* out, int n)  { unsigned i = threadIdx.x; if (i < (unsigned)n) SIM_ST(out, i, 1); }
36
37// initcheck: read a buffer that was never written vs. one that was.
38__global__ void kReadInto(const int* in, int* out) { unsigned i = threadIdx.x; SIM_ST(out, i, SIM_LD(in, i)); }
39
40// racecheck: many threads write the same address with no sync vs. distinct addresses.
41__global__ void kRaceWrite(int* out) { SIM_ST(out, 0, (int)threadIdx.x); }
42__global__ void kNoRace(int* out)    { unsigned i = threadIdx.x; SIM_ST(out, i, (int)i); }
43
44// synccheck: __syncthreads reached by only some threads vs. all.
45__global__ void kSyncDiverge(int* out) {
46    unsigned t = threadIdx.x;
47    if (t < 16) { SIM_ST(out, t, 1); return; }  // these exit before the barrier
48    __syncthreads();                             // only t >= 16 arrive -> divergence
49    SIM_ST(out, t, 2);
SIM_ST (test_sanitizer.cu:34)
24            std::fprintf(stderr, "CUDA error %s at %s:%d: %s\n",             \
25                         cudaGetErrorName(err_), __FILE__, __LINE__,         \
26                         cudaGetErrorString(err_));                          \
27            std::exit(EXIT_FAILURE);                                         \
28        }                                                                    \
29    } while (0)
30
31constexpr int BLK = 32;
32
33// memcheck: out-of-bounds store vs. in-bounds store.
34__global__ void kOOB(int* out, int n)      { SIM_ST(out, threadIdx.x + n, 1); }
35__global__ void kInBounds(int* out, int n)  { unsigned i = threadIdx.x; if (i < (unsigned)n) SIM_ST(out, i, 1); }
36
37// initcheck: read a buffer that was never written vs. one that was.
38__global__ void kReadInto(const int* in, int* out) { unsigned i = threadIdx.x; SIM_ST(out, i, SIM_LD(in, i)); }
39
40// racecheck: many threads write the same address with no sync vs. distinct addresses.
41__global__ void kRaceWrite(int* out) { SIM_ST(out, 0, (int)threadIdx.x); }
42__global__ void kNoRace(int* out)    { unsigned i = threadIdx.x; SIM_ST(out, i, (int)i); }
43
44// synccheck: __syncthreads reached by only some threads vs. all.
45__global__ void kSyncDiverge(int* out) {
simCheckReset (test_sanitizer.cu:76)
66    int deviceCount = 0;
67    if (cudaGetDeviceCount(&deviceCount) != cudaSuccess || deviceCount == 0) { std::printf("no device\n"); return 0; }
68    cudaDeviceProp prop{}; CUDA_CHECK(cudaGetDeviceProperties(&prop, 0));
69    std::printf("Device 0: %s (compute %d.%d)\n\n", prop.name, prop.major, prop.minor);
70
71    bool all = true;
72    int* out = nullptr; CUDA_CHECK(cudaMalloc(&out, BLK * sizeof(int)));
73    int* in  = nullptr; CUDA_CHECK(cudaMalloc(&in,  BLK * sizeof(int)));
74
75    // ---- memcheck --------------------------------------------------------
76    simCheckReset();
77    LAUNCH(kOOB, 1, BLK, out, BLK);
78    all &= expect("memcheck: OOB store", simMemcheckErrors() > 0, true);
79    simCheckReset();
80    LAUNCH(kInBounds, 1, BLK, out, BLK);
81    all &= expect("memcheck: in-bounds store", simMemcheckErrors() > 0, false);
82
83    // ---- initcheck -------------------------------------------------------
84    simCheckReset();
85    LAUNCH(kReadInto, 1, BLK, in, out);          // `in` never initialised
86    all &= expect("initcheck: read uninitialised", simInitcheckErrors() > 0, true);
87    simCheckReset();
simInitcheckErrors (test_sanitizer.cu:86)
76    simCheckReset();
77    LAUNCH(kOOB, 1, BLK, out, BLK);
78    all &= expect("memcheck: OOB store", simMemcheckErrors() > 0, true);
79    simCheckReset();
80    LAUNCH(kInBounds, 1, BLK, out, BLK);
81    all &= expect("memcheck: in-bounds store", simMemcheckErrors() > 0, false);
82
83    // ---- initcheck -------------------------------------------------------
84    simCheckReset();
85    LAUNCH(kReadInto, 1, BLK, in, out);          // `in` never initialised
86    all &= expect("initcheck: read uninitialised", simInitcheckErrors() > 0, true);
87    simCheckReset();
88    CUDA_CHECK(cudaMemset(in, 0, BLK * sizeof(int)));  // now initialised
89    LAUNCH(kReadInto, 1, BLK, in, out);
90    all &= expect("initcheck: read initialised", simInitcheckErrors() > 0, false);
91
92    // ---- racecheck -------------------------------------------------------
93    simCheckReset();
94    LAUNCH(kRaceWrite, 1, BLK, out);             // all write address 0, no sync
95    all &= expect("racecheck: unsynced same addr", simRacecheckErrors() > 0, true);
96    simCheckReset();
97    LAUNCH(kNoRace, 1, BLK, out);                // distinct addresses
simMemcheckErrors (test_sanitizer.cu:78)
68    cudaDeviceProp prop{}; CUDA_CHECK(cudaGetDeviceProperties(&prop, 0));
69    std::printf("Device 0: %s (compute %d.%d)\n\n", prop.name, prop.major, prop.minor);
70
71    bool all = true;
72    int* out = nullptr; CUDA_CHECK(cudaMalloc(&out, BLK * sizeof(int)));
73    int* in  = nullptr; CUDA_CHECK(cudaMalloc(&in,  BLK * sizeof(int)));
74
75    // ---- memcheck --------------------------------------------------------
76    simCheckReset();
77    LAUNCH(kOOB, 1, BLK, out, BLK);
78    all &= expect("memcheck: OOB store", simMemcheckErrors() > 0, true);
79    simCheckReset();
80    LAUNCH(kInBounds, 1, BLK, out, BLK);
81    all &= expect("memcheck: in-bounds store", simMemcheckErrors() > 0, false);
82
83    // ---- initcheck -------------------------------------------------------
84    simCheckReset();
85    LAUNCH(kReadInto, 1, BLK, in, out);          // `in` never initialised
86    all &= expect("initcheck: read uninitialised", simInitcheckErrors() > 0, true);
87    simCheckReset();
88    CUDA_CHECK(cudaMemset(in, 0, BLK * sizeof(int)));  // now initialised
89    LAUNCH(kReadInto, 1, BLK, in, out);
simRacecheckErrors (test_sanitizer.cu:95)
 85    LAUNCH(kReadInto, 1, BLK, in, out);          // `in` never initialised
 86    all &= expect("initcheck: read uninitialised", simInitcheckErrors() > 0, true);
 87    simCheckReset();
 88    CUDA_CHECK(cudaMemset(in, 0, BLK * sizeof(int)));  // now initialised
 89    LAUNCH(kReadInto, 1, BLK, in, out);
 90    all &= expect("initcheck: read initialised", simInitcheckErrors() > 0, false);
 91
 92    // ---- racecheck -------------------------------------------------------
 93    simCheckReset();
 94    LAUNCH(kRaceWrite, 1, BLK, out);             // all write address 0, no sync
 95    all &= expect("racecheck: unsynced same addr", simRacecheckErrors() > 0, true);
 96    simCheckReset();
 97    LAUNCH(kNoRace, 1, BLK, out);                // distinct addresses
 98    all &= expect("racecheck: distinct addrs", simRacecheckErrors() > 0, false);
 99
100    // ---- synccheck -------------------------------------------------------
101    simCheckReset();
102    LAUNCH(kSyncDiverge, 1, BLK, out);
103    all &= expect("synccheck: divergent barrier", simSynccheckErrors() > 0, true);
104    simCheckReset();
105    LAUNCH(kSyncClean, 1, BLK, out);
106    all &= expect("synccheck: uniform barrier", simSynccheckErrors() > 0, false);
simSynccheckErrors (test_sanitizer.cu:103)
 93    simCheckReset();
 94    LAUNCH(kRaceWrite, 1, BLK, out);             // all write address 0, no sync
 95    all &= expect("racecheck: unsynced same addr", simRacecheckErrors() > 0, true);
 96    simCheckReset();
 97    LAUNCH(kNoRace, 1, BLK, out);                // distinct addresses
 98    all &= expect("racecheck: distinct addrs", simRacecheckErrors() > 0, false);
 99
100    // ---- synccheck -------------------------------------------------------
101    simCheckReset();
102    LAUNCH(kSyncDiverge, 1, BLK, out);
103    all &= expect("synccheck: divergent barrier", simSynccheckErrors() > 0, true);
104    simCheckReset();
105    LAUNCH(kSyncClean, 1, BLK, out);
106    all &= expect("synccheck: uniform barrier", simSynccheckErrors() > 0, false);
107
108    CUDA_CHECK(cudaFree(out));
109    CUDA_CHECK(cudaFree(in));
110
111    std::printf("\ntest_sanitizer: %s\n", all ? "ALL PASSED" : "FAILED");
112    return all ? 0 : 1;
113}
On this page