cudaMemset / managed memory#
Set count bytes to value (as unsigned char), like std::memset.
Defined in code/simulator.h.
This page documents 5 symbols: 5 functions.
Functions#
Signature |
Description |
Availability |
Location |
Example |
|---|---|---|---|---|
|
sim only |
simulator.h:2115 |
||
|
sim only |
simulator.h:2110 |
||
|
Pinned (page-locked) host memory: one address space in the sim, so there is no real DMA to accelerate — a page-locked host allocation is just a plain malloc. |
sim only |
simulator.h:2105 |
|
|
Unified memory: one address space in the sim, so managed == a tracked malloc that is directly readable/writable from host code too. |
sim only |
simulator.h:2098 |
|
|
Set |
sim only |
simulator.h:2090 |
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.
cudaFreeHost (test_web_bf_stream_kernel_transfer.cu:109)
99 // Self-check: h_z must equal x + y everywhere.
100 bool ok = true;
101 for (int n = 0; n < N; ++n)
102 if (std::fabs(h_z[n] - (h_x[n] + h_y[n])) > 1e-4) { ok = false; break; }
103 std::printf("kernel-transfer overlap (N=%d, %d streams): %s\n",
104 N, num, ok ? "PASSED" : "FAILED");
105
106 for (int i = 0; i < MAX_NUM_STREAMS; i++)
107 CHECK(cudaStreamDestroy(streams[i]));
108 CHECK(cudaFreeHost(h_x));
109 CHECK(cudaFreeHost(h_y));
110 CHECK(cudaFreeHost(h_z));
111 CHECK(cudaFree(d_x));
112 CHECK(cudaFree(d_y));
113 CHECK(cudaFree(d_z));
114 return ok ? 0 : 1;
115}
cudaHostAlloc (test_web_cs_multicopy.cu:100)
90 memsize = N * sizeof(int);
91
92 thread_blocks = N / block.x;
93 grid.x = thread_blocks % 65535;
94 grid.y = (thread_blocks / 65535 + 1);
95
96 // Allocate resources
97 h_data_source.assign(N, 0); // [SIM] std::vector, all zeros
98
99 for (int i = 0; i < STREAM_COUNT; ++i) {
100 checkCudaErrors(cudaHostAlloc(&h_data_in[i], memsize, cudaHostAllocDefault));
101 checkCudaErrors(cudaMalloc(&d_data_in[i], memsize));
102 checkCudaErrors(cudaMemset(d_data_in[i], 0, memsize));
103
104 checkCudaErrors(cudaHostAlloc(&h_data_out[i], memsize, cudaHostAllocDefault));
105 checkCudaErrors(cudaMalloc(&d_data_out[i], memsize));
106
107 checkCudaErrors(cudaStreamCreate(&stream[i]));
108 checkCudaErrors(cudaEventCreate(&cycleDone[i]));
109
110 cudaEventRecord(cycleDone[i], stream[i]);
111 }
cudaMallocHost (test_web_bf_stream_kernel_transfer.cu:64)
54 if (e != cudaSuccess || deviceCount == 0) {
55 std::printf("No CUDA device available (error %d: %s).\n"
56 "Compilation and toolkit are fine; kernels need an NVIDIA GPU.\n",
57 static_cast<int>(e), cudaGetErrorString(e));
58 return 0;
59 }
60 cudaDeviceProp prop{}; cudaGetDeviceProperties(&prop, 0);
61 std::printf("Device 0: %s (compute %d.%d)\n", prop.name, prop.major, prop.minor);
62
63 real *h_x, *h_y, *h_z;
64 CHECK(cudaMallocHost(&h_x, M)); // pinned host memory (loop-filled below)
65 CHECK(cudaMallocHost(&h_y, M));
66 CHECK(cudaMallocHost(&h_z, M));
67 for (int n = 0; n < N; ++n) { h_x[n] = (real)1.23; h_y[n] = (real)2.34; }
68
69 real *d_x, *d_y, *d_z;
70 CHECK(cudaMalloc(&d_x, M));
71 CHECK(cudaMalloc(&d_y, M));
72 CHECK(cudaMalloc(&d_z, M));
73
74 for (int i = 0; i < MAX_NUM_STREAMS; i++)
75 CHECK(cudaStreamCreate(&(streams[i])));
cudaMallocManaged (test_runtime_api.cu:65)
55 cudaDeviceProp prop{};
56 CUDA_CHECK(cudaGetDeviceProperties(&prop, 0));
57 std::printf("Device 0: %s (compute %d.%d)\n\n", prop.name, prop.major, prop.minor);
58
59 bool all = true;
60 const int blocks = N / BLOCK;
61
62 // Managed (unified) memory: host code reads/writes the same pointers.
63 float *a = nullptr, *b = nullptr, *c = nullptr;
64 CUDA_CHECK(cudaMallocManaged(&a, N * sizeof(float)));
65 CUDA_CHECK(cudaMallocManaged(&b, N * sizeof(float)));
66 CUDA_CHECK(cudaMallocManaged(&c, N * sizeof(float)));
67 for (int i = 0; i < N; ++i) { a[i] = static_cast<float>(i); b[i] = static_cast<float>(2 * i); }
68
69 // cudaMemset must zero the buffer.
70 CUDA_CHECK(cudaMemset(c, 0, N * sizeof(float)));
71 { bool z = true; for (int i = 0; i < N; ++i) if (c[i] != 0.0f) z = false;
72 all &= report("cudaMemset zeroes buffer", z); }
73
74 // Stream + events around a launch.
75 cudaStream_t stream; CUDA_CHECK(cudaStreamCreate(&stream));
cudaMemset (test_debug.cu:51)
41 for (int s = 0; s < NSEG; ++s) {
42 rank[s * BLOCK + static_cast<int>(tid)] = atomicAdd(&ctr[s], 1);
43 __syncthreads();
44 }
45}
46
47static std::vector<int> run() {
48 int *dRank = nullptr, *dCtr = nullptr;
49 CUDA_CHECK(cudaMalloc(&dRank, NSEG * BLOCK * sizeof(int)));
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;