Step-debugger hook (reproducible mid-launch inspection)#

Install a callback the scheduler runs at the start of every block and after each __syncthreads() release. It gets a SimDebug handle to read which block / barrier-segment is executing and — the point of it — to CHANGE the thread visit order for the UPCOMING segment. So an agent can “break at the 2nd reduction, reorder the threads, continue one iteration” deterministically, then combine it with a printf/SIM_LD in the kernel to display values (e.g. tid). CPU-sim only; a no-op under nvcc (real GPUs schedule in hardware).

Defined in code/simulator.h.

This page documents 12 symbols: 1 classes, 11 functions.

Classes#

Signature

Description

Availability

Location

Example

class SimDebug

Step-debugger hook — CPU-sim only; stubbed so a debugging harness compiles and runs on a real GPU (the hook simply never fires there).

sim + nvcc

simulator.h:694

Functions#

Signature

Description

Availability

Location

Example

unsigned block() const

blockIdx.x of the running block.

sim + nvcc

simulator.h:698

View

std::vector<char> seen(static_cast<size_t>(n < 0 ? 0 : n), 0)

sim only

simulator.h:713

View

int segment() const

__syncthreads() barriers passed in this block so far: 0 = block start (before the first barrier), 1 = after the first, 2 = after the second, …

sim + nvcc

simulator.h:701

View

void setThreadOrder(const std::vector<int>& order)

Pin the visit order for the upcoming segment(s) of THIS block.

sim + nvcc

simulator.h:711

View

void setThreadOrder(std::initializer_list<int> order)

sim + nvcc

simulator.h:722

View

inline void simClearStepHook()

sim + nvcc

simulator.h:738

View

explicit SimDebug(simdetail::StepState* s) : s_(s)

sim only

simulator.h:696

inline void simSetStepHook(simdetail::StepHook hook)

Install / remove the step hook (call before a LAUNCH).

sim + nvcc

simulator.h:737

View

int threadCount() const

sim + nvcc

simulator.h:702

std::vector<int> threadOrder() const

The visit order the upcoming segment will use (linear thread ids).

sim + nvcc

simulator.h:704

static void warnBad(size_t sz)

sim only

simulator.h:724

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.

block (test_debug.cu:97)
 87    // The hook: break at the 2nd reduction (segment 2), print the block/segment,
 88    // and reorder the threads for that reduction only.
 89    const std::vector<int> custom = {3, 0, 1, 2, 7, 4, 5, 6};
 90    std::vector<int> segmentsSeen;
 91
 92    simSetThreadOrder(OrderThread::sequential());   // default order for all segments
 93    simSetStepHook([&](SimDebug& d) {
 94        segmentsSeen.push_back(d.segment());
 95        if (d.segment() == 2) {
 96            std::printf("  [hook] break at block %u, segment %d (2nd reduction): "
 97                        "reorder threads -> {3,0,1,2,7,4,5,6}\n", d.block(), d.segment());
 98            d.setThreadOrder(custom);
 99        }
100    });
101
102    std::vector<int> rank = run();
103    simClearStepHook();
104
105    // Show the recorded visit order of each reduction.
106    std::printf("\n  visit ranks per reduction (rank[reduction][tid]):\n");
107    for (int s = 0; s < NSEG; ++s) {
108        std::printf("    reduction %d:", s + 1);
seen (test_debug.cu:60)
50    CUDA_CHECK(cudaMalloc(&dCtr, NSEG * sizeof(int)));
51    CUDA_CHECK(cudaMemset(dCtr, 0, NSEG * sizeof(int)));
52    LAUNCH(kReductions, 1, BLOCK, dRank, dCtr);
53    std::vector<int> rank(NSEG * BLOCK);
54    CUDA_CHECK(cudaMemcpy(rank.data(), dRank, rank.size() * sizeof(int), cudaMemcpyDeviceToHost));
55    CUDA_CHECK(cudaFree(dRank)); CUDA_CHECK(cudaFree(dCtr));
56    return rank;
57}
58
59static bool isPermutation(const int* v, int n) {
60    std::vector<char> seen(n, 0);
61    for (int i = 0; i < n; ++i) { int x = v[i]; if (x < 0 || x >= n || seen[x]) return false; seen[x] = 1; }
62    return true;
63}
64static bool matchesOrder(const int* rank, const std::vector<int>& perm) {
65    for (int k = 0; k < static_cast<int>(perm.size()); ++k) if (rank[perm[k]] != k) return false;
66    return true;
67}
68
69int main() {
70    int deviceCount = 0;
71    cudaError_t err = cudaGetDeviceCount(&deviceCount);
segment (test_debug.cu:94)
 84        true;
 85#endif
 86
 87    // The hook: break at the 2nd reduction (segment 2), print the block/segment,
 88    // and reorder the threads for that reduction only.
 89    const std::vector<int> custom = {3, 0, 1, 2, 7, 4, 5, 6};
 90    std::vector<int> segmentsSeen;
 91
 92    simSetThreadOrder(OrderThread::sequential());   // default order for all segments
 93    simSetStepHook([&](SimDebug& d) {
 94        segmentsSeen.push_back(d.segment());
 95        if (d.segment() == 2) {
 96            std::printf("  [hook] break at block %u, segment %d (2nd reduction): "
 97                        "reorder threads -> {3,0,1,2,7,4,5,6}\n", d.block(), d.segment());
 98            d.setThreadOrder(custom);
 99        }
100    });
101
102    std::vector<int> rank = run();
103    simClearStepHook();
104
105    // Show the recorded visit order of each reduction.
setThreadOrder (test_debug.cu:98)
 88    // and reorder the threads for that reduction only.
 89    const std::vector<int> custom = {3, 0, 1, 2, 7, 4, 5, 6};
 90    std::vector<int> segmentsSeen;
 91
 92    simSetThreadOrder(OrderThread::sequential());   // default order for all segments
 93    simSetStepHook([&](SimDebug& d) {
 94        segmentsSeen.push_back(d.segment());
 95        if (d.segment() == 2) {
 96            std::printf("  [hook] break at block %u, segment %d (2nd reduction): "
 97                        "reorder threads -> {3,0,1,2,7,4,5,6}\n", d.block(), d.segment());
 98            d.setThreadOrder(custom);
 99        }
100    });
101
102    std::vector<int> rank = run();
103    simClearStepHook();
104
105    // Show the recorded visit order of each reduction.
106    std::printf("\n  visit ranks per reduction (rank[reduction][tid]):\n");
107    for (int s = 0; s < NSEG; ++s) {
108        std::printf("    reduction %d:", s + 1);
109        for (int t = 0; t < BLOCK; ++t) std::printf(" %d", rank[s * BLOCK + t]);
setThreadOrder (test_debug.cu:98)
 88    // and reorder the threads for that reduction only.
 89    const std::vector<int> custom = {3, 0, 1, 2, 7, 4, 5, 6};
 90    std::vector<int> segmentsSeen;
 91
 92    simSetThreadOrder(OrderThread::sequential());   // default order for all segments
 93    simSetStepHook([&](SimDebug& d) {
 94        segmentsSeen.push_back(d.segment());
 95        if (d.segment() == 2) {
 96            std::printf("  [hook] break at block %u, segment %d (2nd reduction): "
 97                        "reorder threads -> {3,0,1,2,7,4,5,6}\n", d.block(), d.segment());
 98            d.setThreadOrder(custom);
 99        }
100    });
101
102    std::vector<int> rank = run();
103    simClearStepHook();
104
105    // Show the recorded visit order of each reduction.
106    std::printf("\n  visit ranks per reduction (rank[reduction][tid]):\n");
107    for (int s = 0; s < NSEG; ++s) {
108        std::printf("    reduction %d:", s + 1);
109        for (int t = 0; t < BLOCK; ++t) std::printf(" %d", rank[s * BLOCK + t]);
simClearStepHook (test_debug.cu:103)
 93    simSetStepHook([&](SimDebug& d) {
 94        segmentsSeen.push_back(d.segment());
 95        if (d.segment() == 2) {
 96            std::printf("  [hook] break at block %u, segment %d (2nd reduction): "
 97                        "reorder threads -> {3,0,1,2,7,4,5,6}\n", d.block(), d.segment());
 98            d.setThreadOrder(custom);
 99        }
100    });
101
102    std::vector<int> rank = run();
103    simClearStepHook();
104
105    // Show the recorded visit order of each reduction.
106    std::printf("\n  visit ranks per reduction (rank[reduction][tid]):\n");
107    for (int s = 0; s < NSEG; ++s) {
108        std::printf("    reduction %d:", s + 1);
109        for (int t = 0; t < BLOCK; ++t) std::printf(" %d", rank[s * BLOCK + t]);
110        std::printf("\n");
111    }
112    std::printf("  hook fired at segments:");
113    for (int s : segmentsSeen) std::printf(" %d", s);
114    std::printf("\n\n");
simSetStepHook (test_debug.cu:93)
 83#else
 84        true;
 85#endif
 86
 87    // The hook: break at the 2nd reduction (segment 2), print the block/segment,
 88    // and reorder the threads for that reduction only.
 89    const std::vector<int> custom = {3, 0, 1, 2, 7, 4, 5, 6};
 90    std::vector<int> segmentsSeen;
 91
 92    simSetThreadOrder(OrderThread::sequential());   // default order for all segments
 93    simSetStepHook([&](SimDebug& d) {
 94        segmentsSeen.push_back(d.segment());
 95        if (d.segment() == 2) {
 96            std::printf("  [hook] break at block %u, segment %d (2nd reduction): "
 97                        "reorder threads -> {3,0,1,2,7,4,5,6}\n", d.block(), d.segment());
 98            d.setThreadOrder(custom);
 99        }
100    });
101
102    std::vector<int> rank = run();
103    simClearStepHook();
On this page