Device query / error stubs#
The sim reports a single device that mimics the sm_75 target (compute 7.5, 2 GiB), and — since every operation is already synchronous and cannot fail asynchronously — the sync/error queries always report success.
Defined in code/simulator.h.
This page documents 10 symbols: 10 functions.
Functions#
Signature |
Description |
Availability |
Location |
Example |
|---|---|---|---|---|
|
Context reset is a no-op in the sim. |
sim only |
simulator.h:2063 |
|
|
No-op: launches are synchronous in the sim, so there is nothing to wait for. |
sim only |
simulator.h:2055 |
|
|
sim only |
simulator.h:2061 |
||
|
Always reports exactly one (simulated) device. |
sim only |
simulator.h:2014 |
|
|
Fills in a plausible device 0: “CUDA Simulator (CPU, fibers)”, compute 7.5, 2 GiB. |
sim only |
simulator.h:2016 |
|
|
Map an error code to its enum spelling / human-readable message. |
sim only |
simulator.h:2065 |
|
|
sim only |
simulator.h:2076 |
||
|
Error queries: launches cannot fail here, so both always report success. |
sim only |
simulator.h:2057 |
|
|
sim only |
simulator.h:2058 |
||
|
Device selection is a no-op (only device 0 exists). |
sim only |
simulator.h:2060 |
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.
cudaDeviceReset (test_web_enccs_matidx.cu:99)
89 // CPU-reference self-check: the 2-D->linear mapping must cover every idx in
90 // [0,nxy) exactly once, and each thread must have read value == its idx.
91 bool ok = true;
92 for (int i = 0; i < nxy; i++)
93 if (h_matrixA[i] != i || h_out[i] != i) { ok = false; break; }
94
95 std::printf("\n2-D thread-index mapping (%dx%d): %s\n", nx, ny, ok ? "PASSED" : "FAILED");
96
97 cudaFree(d_matrixA); cudaFree(d_out);
98 CHECK(cudaDeviceReset());
99 return ok ? 0 : 1;
100}
cudaDeviceSynchronize (test_indexing.cu:80)
70 std::vector<int> host(static_cast<size_t>(total), SENTINEL);
71
72 int* dev = nullptr;
73 CUDA_CHECK(cudaMalloc(&dev, static_cast<size_t>(total) * sizeof(int)));
74 CUDA_CHECK(cudaMemcpy(dev, host.data(), static_cast<size_t>(total) * sizeof(int),
75 cudaMemcpyHostToDevice));
76
77 LAUNCH(writeGid, grid, block, dev, total);
78 CUDA_CHECK(cudaGetLastError());
79 CUDA_CHECK(cudaDeviceSynchronize());
80
81 CUDA_CHECK(cudaMemcpy(host.data(), dev, static_cast<size_t>(total) * sizeof(int),
82 cudaMemcpyDeviceToHost));
83 CUDA_CHECK(cudaFree(dev));
84
85 // Every cell must equal its index: no gaps (missing thread) and no poison
86 // (out-of-range built-in).
87 int firstBad = -1;
88 for (int i = 0; i < total; ++i) {
89 if (host[static_cast<size_t>(i)] != i) { firstBad = i; break; }
90 }
cudaGetDevice (test_runtime_api.cu:110)
100 CUDA_CHECK(cudaStreamSynchronize(stream));
101 { bool ok = true; for (int i = 0; i < N; ++i) if (back[i] != c[i]) ok = false;
102 all &= report("async memset + memcpy round-trip", ok); }
103
104 // Error-code round-trip through the widened table.
105 {
106 bool ok = true;
107 ok &= (std::string(cudaGetErrorName(cudaErrorNotReady)) == "cudaErrorNotReady");
108 ok &= (std::string(cudaGetErrorString(cudaErrorInvalidDevicePointer)) == "invalid device pointer");
109 ok &= (cudaPeekAtLastError() == cudaSuccess);
110 int dev = -1; ok &= (cudaGetDevice(&dev) == cudaSuccess) && dev == 0;
111 all &= report("error/query round-trip", ok);
112 }
113
114 CUDA_CHECK(cudaEventDestroy(startE));
115 CUDA_CHECK(cudaEventDestroy(stopE));
116 CUDA_CHECK(cudaStreamDestroy(stream));
117 CUDA_CHECK(cudaFree(d));
118 CUDA_CHECK(cudaFree(a));
119 CUDA_CHECK(cudaFree(b));
120 CUDA_CHECK(cudaFree(c));
cudaGetDeviceCount (test_atomics.cu:208)
198 // atomicSub countdown from N -> 0.
199 { int* d = devFrom<int>(N); LAUNCH(kSub, BLOCKS, BLOCK, d);
200 ok &= report("atomicSub countdown", devRead(d) == 0, seed); }
201
202 return ok;
203}
204
205int main() {
206 int deviceCount = 0;
207 cudaError_t err = cudaGetDeviceCount(&deviceCount);
208 if (err != cudaSuccess || deviceCount == 0) {
209 const char* reason = cudaGetErrorString(err);
210 std::printf("No CUDA device available (error %d: %s).\n"
211 "Compilation and toolkit are fine; kernels need an NVIDIA GPU.\n",
212 static_cast<int>(err), reason ? reason : "unknown");
213 return 0;
214 }
215
216 cudaDeviceProp prop{};
217 CUDA_CHECK(cudaGetDeviceProperties(&prop, 0));
218 std::printf("Device 0: %s (compute %d.%d)\n\n", prop.name, prop.major, prop.minor);
cudaGetDeviceProperties (test_atomics.cu:218)
208 cudaError_t err = cudaGetDeviceCount(&deviceCount);
209 if (err != cudaSuccess || deviceCount == 0) {
210 const char* reason = cudaGetErrorString(err);
211 std::printf("No CUDA device available (error %d: %s).\n"
212 "Compilation and toolkit are fine; kernels need an NVIDIA GPU.\n",
213 static_cast<int>(err), reason ? reason : "unknown");
214 return 0;
215 }
216
217 cudaDeviceProp prop{};
218 CUDA_CHECK(cudaGetDeviceProperties(&prop, 0));
219 std::printf("Device 0: %s (compute %d.%d)\n\n", prop.name, prop.major, prop.minor);
220
221 bool all = true;
222 for (unsigned seed = 1; seed <= 4; ++seed) {
223 std::printf("--- schedule seed %u ---\n", seed);
224 all &= runAll(seed);
225 }
226
227 std::printf("\ntest_atomics: %s\n", all ? "ALL PASSED" : "FAILED");
228 return all ? 0 : 1;
229}
cudaGetErrorName (test_atomics.cu:24)
14#include <cstdio>
15#include <cstdlib>
16#include <vector>
17
18#define CUDA_CHECK(call) \
19 do { \
20 cudaError_t err_ = (call); \
21 if (err_ != cudaSuccess) { \
22 std::fprintf(stderr, "CUDA error %s at %s:%d: %s\n", \
23 cudaGetErrorName(err_), __FILE__, __LINE__, \
24 cudaGetErrorString(err_)); \
25 std::exit(EXIT_FAILURE); \
26 } \
27 } while (0)
28
29constexpr int N = 4096; // number of contributing threads
30constexpr int BLOCK = 128;
31constexpr int BLOCKS = N / BLOCK;
32constexpr int BINS = 16;
33
34// ---- Kernels ---------------------------------------------------------------
cudaGetErrorString (test_atomics.cu:25)
15#include <cstdio>
16#include <cstdlib>
17#include <vector>
18
19#define CUDA_CHECK(call) \
20 do { \
21 cudaError_t err_ = (call); \
22 if (err_ != cudaSuccess) { \
23 std::fprintf(stderr, "CUDA error %s at %s:%d: %s\n", \
24 cudaGetErrorName(err_), __FILE__, __LINE__, \
25 cudaGetErrorString(err_)); \
26 std::exit(EXIT_FAILURE); \
27 } \
28 } while (0)
29
30constexpr int N = 4096; // number of contributing threads
31constexpr int BLOCK = 128;
32constexpr int BLOCKS = N / BLOCK;
33constexpr int BINS = 16;
34
35// ---- Kernels ---------------------------------------------------------------
cudaGetLastError (test_indexing.cu:79)
69 static_cast<int>(block.x * block.y * block.z);
70
71 std::vector<int> host(static_cast<size_t>(total), SENTINEL);
72
73 int* dev = nullptr;
74 CUDA_CHECK(cudaMalloc(&dev, static_cast<size_t>(total) * sizeof(int)));
75 CUDA_CHECK(cudaMemcpy(dev, host.data(), static_cast<size_t>(total) * sizeof(int),
76 cudaMemcpyHostToDevice));
77
78 LAUNCH(writeGid, grid, block, dev, total);
79 CUDA_CHECK(cudaGetLastError());
80 CUDA_CHECK(cudaDeviceSynchronize());
81
82 CUDA_CHECK(cudaMemcpy(host.data(), dev, static_cast<size_t>(total) * sizeof(int),
83 cudaMemcpyDeviceToHost));
84 CUDA_CHECK(cudaFree(dev));
85
86 // Every cell must equal its index: no gaps (missing thread) and no poison
87 // (out-of-range built-in).
88 int firstBad = -1;
89 for (int i = 0; i < total; ++i) {
90 if (host[static_cast<size_t>(i)] != i) { firstBad = i; break; }
cudaPeekAtLastError (test_runtime_api.cu:109)
99 CUDA_CHECK(cudaMemcpyAsync(back.data(), d, N * sizeof(float), cudaMemcpyDeviceToHost, stream));
100 CUDA_CHECK(cudaStreamSynchronize(stream));
101 { bool ok = true; for (int i = 0; i < N; ++i) if (back[i] != c[i]) ok = false;
102 all &= report("async memset + memcpy round-trip", ok); }
103
104 // Error-code round-trip through the widened table.
105 {
106 bool ok = true;
107 ok &= (std::string(cudaGetErrorName(cudaErrorNotReady)) == "cudaErrorNotReady");
108 ok &= (std::string(cudaGetErrorString(cudaErrorInvalidDevicePointer)) == "invalid device pointer");
109 ok &= (cudaPeekAtLastError() == cudaSuccess);
110 int dev = -1; ok &= (cudaGetDevice(&dev) == cudaSuccess) && dev == 0;
111 all &= report("error/query round-trip", ok);
112 }
113
114 CUDA_CHECK(cudaEventDestroy(startE));
115 CUDA_CHECK(cudaEventDestroy(stopE));
116 CUDA_CHECK(cudaStreamDestroy(stream));
117 CUDA_CHECK(cudaFree(d));
118 CUDA_CHECK(cudaFree(a));
119 CUDA_CHECK(cudaFree(b));
120 CUDA_CHECK(cudaFree(c));
cudaSetDevice (test_web_bf_query.cu:37)
27 int deviceCount = 0;
28 cudaError_t e = cudaGetDeviceCount(&deviceCount);
29 if (e != cudaSuccess || deviceCount == 0) {
30 std::printf("No CUDA device available (error %d: %s).\n"
31 "Compilation and toolkit are fine; kernels need an NVIDIA GPU.\n",
32 static_cast<int>(e), cudaGetErrorString(e));
33 return 0;
34 }
35
36 CHECK(cudaSetDevice(device_id));
37
38 cudaDeviceProp prop{};
39 CHECK(cudaGetDeviceProperties(&prop, device_id));
40
41 // [SIM] only the fields the sim's cudaDeviceProp carries.
42 std::printf("Device id: %d\n", device_id);
43 std::printf("Device name: %s\n", prop.name);
44 std::printf("Compute capability: %d.%d\n", prop.major, prop.minor);
45 std::printf("Amount of global memory: %g GB\n",
46 prop.totalGlobalMem / (1024.0 * 1024 * 1024));