diff options
Diffstat (limited to 'tools/cuda-ab/streams2.cu')
| -rw-r--r-- | tools/cuda-ab/streams2.cu | 175 |
1 files changed, 175 insertions, 0 deletions
diff --git a/tools/cuda-ab/streams2.cu b/tools/cuda-ab/streams2.cu new file mode 100644 index 0000000..a5dc5c2 --- /dev/null +++ b/tools/cuda-ab/streams2.cu @@ -0,0 +1,175 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * streams2.cu - two non-default streams, concurrent kernels, event ordering. + * + * Two independent vectorAdds run concurrently on two non-default streams. + * A cross-stream dependency is then imposed with events: stream 2's in-place + * scale waits on stream 1's add via cudaEventRecord/cudaStreamWaitEvent, and + * stream 1's copy-back waits on stream 2's scale the same way. Both streams + * are joined with cudaStreamSynchronize before the host verifies: + * + * c1 == a1 + b1 + * c2 == (a2 + b2) * 2 + * + * Output lines are stable and greppable: + * + * DEVICE <name> cc=<major>.<minor> sm=<count> mem=<bytes> + * STREAMS n=<count> bad=<mismatches> + * TIMING ms=<elapsed> + * STATUS OK|FAIL + * + * Build with a static CUDA runtime (--cudart=static) so the guest only needs + * the driver's libcuda.so, not a CUDA toolkit installation: + * + * nvcc -O2 --cudart=static -o streams2 streams2.cu + * + * Then, from the guest: /tmp/streams2 + */ + +#include <cstdio> +#include <cstdlib> +#include <cstdint> +#include <cuda_runtime.h> + +#define CHECK(x) \ + do { \ + cudaError_t e_ = (x); \ + if (e_ != cudaSuccess) { \ + printf("CUDA_ERROR %s at line %d\n", \ + cudaGetErrorString(e_), __LINE__); \ + printf("STATUS FAIL\n"); \ + return 1; \ + } \ + } while (0) + +__global__ void vadd(const float *a, const float *b, float *c, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + c[i] = a[i] + b[i]; +} + +__global__ void scale(float *c, float s, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + c[i] = c[i] * s; +} + +int main(void) +{ + cudaDeviceProp prop; + const int n = 1 << 15; + const float s2 = 2.0f; + size_t bytes = (size_t)n * sizeof(float); + float *h_a1, *h_b1, *h_c1, *h_a2, *h_b2, *h_c2; + float *d_a1, *d_b1, *d_c1, *d_a2, *d_b2, *d_c2; + cudaStream_t s1, s2s; + cudaEvent_t ev_add1, ev_scale2, t0, t1; + float ms = 0.0f; + int bad = 0, i; + const int grid = (n + 255) / 256; + + setvbuf(stdout, NULL, _IONBF, 0); + + CHECK(cudaGetDeviceProperties(&prop, 0)); + printf("DEVICE %s cc=%d.%d sm=%d mem=%llu\n", prop.name, + prop.major, prop.minor, prop.multiProcessorCount, + (unsigned long long)prop.totalGlobalMem); + + h_a1 = (float *)malloc(bytes); + h_b1 = (float *)malloc(bytes); + h_c1 = (float *)malloc(bytes); + h_a2 = (float *)malloc(bytes); + h_b2 = (float *)malloc(bytes); + h_c2 = (float *)malloc(bytes); + if (!h_a1 || !h_b1 || !h_c1 || !h_a2 || !h_b2 || !h_c2) + return 1; + + for (i = 0; i < n; i++) { + h_a1[i] = (float)(i % 1000) * 0.5f; + h_b1[i] = (float)(i % 997) * 0.25f; + h_a2[i] = (float)(i % 991) * 0.125f; + h_b2[i] = (float)(i % 983) * 0.0625f; + } + + CHECK(cudaMalloc(&d_a1, bytes)); + CHECK(cudaMalloc(&d_b1, bytes)); + CHECK(cudaMalloc(&d_c1, bytes)); + CHECK(cudaMalloc(&d_a2, bytes)); + CHECK(cudaMalloc(&d_b2, bytes)); + CHECK(cudaMalloc(&d_c2, bytes)); + + CHECK(cudaStreamCreate(&s1)); + CHECK(cudaStreamCreate(&s2s)); + CHECK(cudaEventCreate(&ev_add1)); + CHECK(cudaEventCreate(&ev_scale2)); + CHECK(cudaEventCreate(&t0)); + CHECK(cudaEventCreate(&t1)); + + CHECK(cudaEventRecord(t0)); + + /* Both input copies and both add kernels run concurrently. */ + CHECK(cudaMemcpyAsync(d_a1, h_a1, bytes, cudaMemcpyHostToDevice, s1)); + CHECK(cudaMemcpyAsync(d_b1, h_b1, bytes, cudaMemcpyHostToDevice, s1)); + CHECK(cudaMemcpyAsync(d_a2, h_a2, bytes, cudaMemcpyHostToDevice, s2s)); + CHECK(cudaMemcpyAsync(d_b2, h_b2, bytes, cudaMemcpyHostToDevice, s2s)); + vadd<<<grid, 256, 0, s1>>>(d_a1, d_b1, d_c1, n); + vadd<<<grid, 256, 0, s2s>>>(d_a2, d_b2, d_c2, n); + CHECK(cudaGetLastError()); + + /* stream 2's scale waits for stream 1's add... */ + CHECK(cudaEventRecord(ev_add1, s1)); + CHECK(cudaStreamWaitEvent(s2s, ev_add1, 0)); + scale<<<grid, 256, 0, s2s>>>(d_c2, s2, n); + CHECK(cudaGetLastError()); + + /* ...and stream 1's copy-back waits for stream 2's scale. */ + CHECK(cudaEventRecord(ev_scale2, s2s)); + CHECK(cudaStreamWaitEvent(s1, ev_scale2, 0)); + CHECK(cudaMemcpyAsync(h_c1, d_c1, bytes, cudaMemcpyDeviceToHost, s1)); + CHECK(cudaMemcpyAsync(h_c2, d_c2, bytes, cudaMemcpyDeviceToHost, s2s)); + CHECK(cudaEventRecord(t1, s1)); + + CHECK(cudaStreamSynchronize(s1)); + CHECK(cudaStreamSynchronize(s2s)); + CHECK(cudaEventElapsedTime(&ms, t0, t1)); + + for (i = 0; i < n; i++) { + if (h_c1[i] != h_a1[i] + h_b1[i]) + bad++; + if (h_c2[i] != (h_a2[i] + h_b2[i]) * s2) + bad++; + } + + printf("STREAMS n=%d bad=%d\n", n, bad); + printf("TIMING ms=%.3f\n", ms); + + if (bad != 0) { + printf("STATUS FAIL\n"); + return 1; + } + + printf("STATUS OK\n"); + CHECK(cudaStreamDestroy(s1)); + CHECK(cudaStreamDestroy(s2s)); + CHECK(cudaEventDestroy(ev_add1)); + CHECK(cudaEventDestroy(ev_scale2)); + CHECK(cudaEventDestroy(t0)); + CHECK(cudaEventDestroy(t1)); + CHECK(cudaFree(d_a1)); + CHECK(cudaFree(d_b1)); + CHECK(cudaFree(d_c1)); + CHECK(cudaFree(d_a2)); + CHECK(cudaFree(d_b2)); + CHECK(cudaFree(d_c2)); + free(h_a1); + free(h_b1); + free(h_c1); + free(h_a2); + free(h_b2); + free(h_c2); + return 0; +} |