Stream-ordered allocation (synchronous here)#

Defined in code/simulator.h.

This page documents 2 symbols: 2 functions.

Functions#

Signature

Description

Availability

Location

Example

inline cudaError_t cudaFreeAsync(void* ptr, cudaStream_t )

sim only

simulator.h:2386

View

inline cudaError_t cudaMallocAsync(T** ptr, size_t size, cudaStream_t )

sim only

simulator.h:2385

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.

cudaFreeAsync (test_web_cis_stream_v2.cu:90)
 80        checkCuda(cudaMemcpyAsync(d_b, &b[offset], streamBytes,
 81                                  cudaMemcpyHostToDevice, stream[i]));
 82
 83        // [SIM][FIX] own buffers + slice-local length (was `*d_a` + global offset).
 84        LAUNCH_STREAM(addKernel, streamSize / blockSize, blockSize, stream[i],
 85                      d_a, d_b, d_c, streamSize);
 86
 87        checkCuda(cudaMemcpyAsync(&c[offset], d_c, streamBytes,
 88                                  cudaMemcpyDeviceToHost, stream[i]));
 89
 90        checkCuda(cudaFreeAsync(d_a, stream[i]));
 91        checkCuda(cudaFreeAsync(d_b, stream[i]));
 92        checkCuda(cudaFreeAsync(d_c, stream[i]));
 93    }
 94
 95    for (int i = 0; i < nStreams; ++i)
 96        checkCuda(cudaStreamDestroy(stream[i]));
 97
 98    // CPU reference: c must equal a + b everywhere.
 99    bool ok = true;
100    for (int i = 0; i < n && ok; ++i)
101        if (std::fabs(c[i] - (a[i] + b[i])) > 1e-3f) ok = false;
cudaMallocAsync (test_web_cis_stream_v2.cu:74)
64    cudaStream_t stream[nStreams];
65    for (int i = 0; i < nStreams; ++i)
66        cudaStreamCreate(&stream[i]);
67
68    for (int i = 0; i < nStreams; ++i) {
69        int offset = i * streamSize;
70        float *d_a = nullptr, *d_b = nullptr, *d_c = nullptr;
71
72        // [SIM][FIX] per-stream buffers via the stream-ordered allocator.
73        checkCuda(cudaMallocAsync((void **)&d_a, streamBytes, stream[i]));
74        checkCuda(cudaMallocAsync((void **)&d_b, streamBytes, stream[i]));
75        checkCuda(cudaMallocAsync((void **)&d_c, streamBytes, stream[i]));
76
77        checkCuda(cudaMemcpyAsync(d_a, &a[offset], streamBytes,
78                                  cudaMemcpyHostToDevice, stream[i]));
79        checkCuda(cudaMemcpyAsync(d_b, &b[offset], streamBytes,
80                                  cudaMemcpyHostToDevice, stream[i]));
81
82        // [SIM][FIX] own buffers + slice-local length (was `*d_a` + global offset).
83        LAUNCH_STREAM(addKernel, streamSize / blockSize, blockSize, stream[i],
84                      d_a, d_b, d_c, streamSize);
On this page