Internal scheduler state#

Defined in code/simulator.h.

This page documents 2 symbols: 2 functions.

Functions#

Signature

Description

Availability

Location

Example

inline void simSetAtomicPreempt(double p)

Probability (0..1) that an atomic operation preempts the calling thread.

sim + nvcc

simulator.h:628

View

inline void simSetSeed(unsigned seed)

Reseed the block-order RNG explicitly (optional; SIM_SEED does the same).

sim + nvcc

simulator.h:621

View

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.

simSetAtomicPreempt (test_web_cbe_hashtable.cu:177)
167                    "Compilation and toolkit are fine; kernels need an NVIDIA GPU.\n",
168                    static_cast<int>(e), cudaGetErrorString(e));
169        return 0;
170    }
171    cudaDeviceProp prop{};
172    cudaGetDeviceProperties(&prop, 0);
173    std::printf("Device 0: %s (compute %d.%d)\n", prop.name, prop.major, prop.minor);
174
175    // [SIM] Preempt at atomics so the per-bucket spin-locks are really
176    // contended. No-op under nvcc — a real GPU interleaves threads anyway.
177    simSetAtomicPreempt(0.5);
178
179    // [SIM] std::vector rather than malloc + fill loop (MSVC /O2 memcpy bug —
180    // see code\test_web_reduce_idle.cu).
181    std::vector<unsigned int> buffer(ELEMENTS);
182    for (size_t i = 0; i < ELEMENTS; i++) buffer[i] = (unsigned int)rand();
183
184    unsigned int* dev_keys   = nullptr;
185    void**        dev_values = nullptr;
186    HANDLE_ERROR(cudaMalloc((void**)&dev_keys, SIZE));
187    HANDLE_ERROR(cudaMalloc((void**)&dev_values, ELEMENTS * sizeof(void*)));
188    HANDLE_ERROR(cudaMemcpy(dev_keys, buffer.data(), SIZE, cudaMemcpyHostToDevice));
simSetSeed (test_atomics.cu:115)
105}
106
107static bool report(const char* name, bool ok, unsigned seed) {
108    std::printf("  %-28s seed=%-5u %s\n", name, seed, ok ? "PASSED" : "FAILED");
109    return ok;
110}
111
112// Runs the full battery once under a given schedule seed.
113static bool runAll(unsigned seed) {
114#ifndef __CUDACC__
115    simSetSeed(seed);
116#endif
117    bool ok = true;
118
119    // Integer / unsigned-long-long / float / double counters -> exactly N.
120    { int* d = devFrom<int>(0); LAUNCH(kCounterInt, BLOCKS, BLOCK, d);
121      ok &= report("atomicAdd int counter", devRead(d) == N, seed); }
122    { unsigned long long* d = devFrom<unsigned long long>(0ULL);
123      LAUNCH(kCounterUll, BLOCKS, BLOCK, d);
124      ok &= report("atomicAdd ull counter", devRead(d) == (unsigned long long)N, seed); }
125    { float* d = devFrom<float>(0.0f); LAUNCH(kCounterFloat, BLOCKS, BLOCK, d);
126      ok &= report("atomicAdd float counter", devRead(d) == (float)N, seed); }
On this page