diff options
Diffstat (limited to 'tools/cuda-ab/p2/src/p2_streams.cu')
| -rw-r--r-- | tools/cuda-ab/p2/src/p2_streams.cu | 86 |
1 files changed, 81 insertions, 5 deletions
diff --git a/tools/cuda-ab/p2/src/p2_streams.cu b/tools/cuda-ab/p2/src/p2_streams.cu index 09d6ed3..12ee56c 100644 --- a/tools/cuda-ab/p2/src/p2_streams.cu +++ b/tools/cuda-ab/p2/src/p2_streams.cu @@ -1,16 +1,23 @@ // SPDX-License-Identifier: BSD-3-Clause /* - * p2_streams - three created streams, two kernels per stream (scale then - * add), independent 1024-float buffers, per-stream events. + * p2_streams - created streams, two kernels per stream, independent buffers, + * per-stream events. * - * Expected value per element: ((i % 100) * 2.0f) + 1.0f (exact in fp32). - * After all streams complete, each event must query as complete. - * Device memory: 3 x 4 KB (3 pages). + * phase 1 three streams, each: scale (x2) then addv (+1), dependent pair + * phase 2 four streams, each: two *independent* fill kernels on separate + * buffers - up to eight kernels in flight simultaneously; all + * results verified and every stream event required complete + * + * Expected value phase 1 per element: ((i % 100) * 2.0f) + 1.0f (exact in + * fp32). Phase 2 values are small integers (exact in int32). + * Device memory: phase 1 3 x 4 KB, phase 2 8 x 4 KB (peak 32 KB, 8 pages). */ #include "common.h" #define SN 1024 #define NS 3 +#define IN_NS 4 +#define IN_N 1024 __global__ void scale(float *x, float s, int n) { @@ -28,6 +35,14 @@ __global__ void addv(float *x, float v, int n) x[i] = x[i] + v; } +__global__ void fillk(int *x, int base, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + x[i] = base * 1000 + i; +} + static void fill(float *h, int n) { for (int i = 0; i < n; i++) @@ -39,6 +54,61 @@ static float expected1(int i) return ((float)(i % 100) * 2.0f) + 1.0f; } +/* Phase 2: independent kernels submitted across four streams. No kernel + * depends on any other, so the driver is free to overlap all eight; every + * buffer is verified element-exactly and every stream event must be complete + * after the stream is synchronized. */ +static int indep_phase(void) +{ + int *got = (int *)malloc(sizeof(int) * IN_N); + cudaStream_t st[IN_NS]; + cudaEvent_t ev[IN_NS]; + int *d_a[IN_NS], *d_b[IN_NS]; + int vbad = 0, ebad = 0; + + if (!got) + return 1; + + for (int k = 0; k < IN_NS; k++) { + P2_CHECK(cudaStreamCreate(&st[k])); + P2_CHECK(cudaEventCreate(&ev[k])); + P2_CHECK(cudaMalloc(&d_a[k], sizeof(int) * IN_N)); + P2_CHECK(cudaMalloc(&d_b[k], sizeof(int) * IN_N)); + fillk<<<(IN_N + 255) / 256, 256, 0, st[k]>>>(d_a[k], k + 2, + IN_N); + fillk<<<(IN_N + 255) / 256, 256, 0, st[k]>>>(d_b[k], k + 9, + IN_N); + P2_CHECK(cudaGetLastError()); + P2_CHECK(cudaEventRecord(ev[k], st[k])); + } + + for (int k = 0; k < IN_NS; k++) { + P2_CHECK(cudaStreamSynchronize(st[k])); + P2_CHECK(cudaMemcpy(got, d_a[k], sizeof(int) * IN_N, + cudaMemcpyDeviceToHost)); + for (int i = 0; i < IN_N; i++) + if (got[i] != (k + 2) * 1000 + i) + vbad++; + P2_CHECK(cudaMemcpy(got, d_b[k], sizeof(int) * IN_N, + cudaMemcpyDeviceToHost)); + for (int i = 0; i < IN_N; i++) + if (got[i] != (k + 9) * 1000 + i) + vbad++; + if (cudaEventQuery(ev[k]) != cudaSuccess) + ebad++; + cudaEventDestroy(ev[k]); + cudaStreamDestroy(st[k]); + cudaFree(d_a[k]); + cudaFree(d_b[k]); + } + + printf("RESULT streams indep nstreams=%d launches=%d n=%d bad=%d\n", + IN_NS, IN_NS * 2, IN_N, vbad); + printf("RESULT streams indep_events=%d bad=%d\n", IN_NS, ebad); + free(got); + return vbad + ebad; +} + int main(int argc, char **argv) { float *h = (float *)malloc(sizeof(float) * SN); @@ -53,6 +123,9 @@ int main(int argc, char **argv) printf("RESULT streams nstreams=%d launches=%d n=%d bad=0\n", NS, NS * 2, SN); printf("RESULT streams events=%d bad=0\n", NS); + printf("RESULT streams indep nstreams=%d launches=%d n=%d bad=0\n", + IN_NS, IN_NS * 2, IN_N); + printf("RESULT streams indep_events=%d bad=0\n", IN_NS); free(h); free(got); return p2_status(0); } @@ -104,6 +177,9 @@ int main(int argc, char **argv) cudaStreamDestroy(st[k]); cudaFree(d_x[k]); } + + bad += indep_phase(); + free(h); free(got); return p2_status(bad); } |