diff options
Diffstat (limited to 'tools/cuda-ab/p2')
| -rw-r--r-- | tools/cuda-ab/p2/README.md | 27 | ||||
| -rw-r--r-- | tools/cuda-ab/p2/expected/SOURCE.txt | 6 | ||||
| -rw-r--r-- | tools/cuda-ab/p2/expected/p2_printf_err.full.txt | 1 | ||||
| -rw-r--r-- | tools/cuda-ab/p2/expected/p2_printf_err.txt | 1 | ||||
| -rw-r--r-- | tools/cuda-ab/p2/expected/p2_regs.full.txt | 5 | ||||
| -rw-r--r-- | tools/cuda-ab/p2/expected/p2_regs.txt | 5 | ||||
| -rw-r--r-- | tools/cuda-ab/p2/expected/p2_streams.full.txt | 2 | ||||
| -rw-r--r-- | tools/cuda-ab/p2/expected/p2_streams.txt | 2 | ||||
| -rwxr-xr-x | tools/cuda-ab/p2/run_guest.sh | 9 | ||||
| -rw-r--r-- | tools/cuda-ab/p2/src/p2_printf_err.cu | 36 | ||||
| -rw-r--r-- | tools/cuda-ab/p2/src/p2_regs.cu | 179 | ||||
| -rw-r--r-- | tools/cuda-ab/p2/src/p2_streams.cu | 86 |
12 files changed, 338 insertions, 21 deletions
diff --git a/tools/cuda-ab/p2/README.md b/tools/cuda-ab/p2/README.md index da4e896..f2e1893 100644 --- a/tools/cuda-ab/p2/README.md +++ b/tools/cuda-ab/p2/README.md @@ -2,17 +2,19 @@ A reusable, self-checking acceptance suite for the Open-IOV emulated vGPU. It extends the single-workload `ab.sh` A/B harness (`victory_mail`, -`vector_add`) to ten independent programs that exercise launch shape, kernel -parameters, shared memory, streams, allocator churn, device printf, error -reporting, a compute-heavy benchmark, and a combined program - each of which -verifies its own results exactly and prints machine-checkable lines. +`vector_add`) to eleven independent programs that exercise launch shape, +kernel parameters, register/launch-bounds variants, shared memory, streams +(including independent simultaneous kernels), allocator churn, device +printf/assert, error reporting, a compute-heavy benchmark, and a combined +program - each of which verifies its own results exactly and prints +machine-checkable lines. ``` p2/ ├── build.sh # nvcc -O2 -arch=sm_75 --cudart=static -> bin/ ├── oracle.sh # host reference run (GPU) or deterministic emit fallback ├── run_guest.sh # push to guest, run with timeout, diff vs expected/, table -├── src/ # common.h + 10 .cu programs +├── src/ # common.h + 11 .cu programs ├── bin/ # built binaries (rebuild with ./build.sh) ├── expected/ # <name>.txt stable lines (+ <name>.full.txt, SOURCE.txt) └── guest/ # per-run guest output (created by run_guest.sh) @@ -25,10 +27,11 @@ p2/ | `p2_seq5` | 5 sequential launches in one process on the same buffers, state verified after each | 3 x 4 KB | | `p2_shapes` | 1x1x1, odd sizes (1013, block 33), non-multiple n (1000/64), 2D grid (40x20), 3D grid (24x8x8) with bounds checks | 8 KB buffer | | `p2_params` | kernel parameter counts 0..12, types int/float/double/device-pointer | 64 B + 8 B + 4 B | +| `p2_regs` | register/launch-bounds variants: `__launch_bounds__` caps, 64-accumulator register pressure, two block sizes; register counts reported as INFO | 4 KB | | `p2_smem` | shared-memory tree reduction (4096 -> 1) and 16x16-tiled 32x32 int GEMM | peak 4 pages | -| `p2_streams` | 3 streams x 2 kernels, per-stream events and result verification | 3 x 4 KB | +| `p2_streams` | 3 streams x 2 dependent kernels, per-stream events; 4 streams x 2 *independent* kernels (8 in flight) with exact checks | 3 x 4 KB, then 8 x 4 KB | | `p2_churn` | 14 cudaMalloc/cudaFree across 6 iterations plus a final fresh alloc | peak 32 KB | -| `p2_printf_err` | device printf (5 deterministic lines) + invalid grid/block launch error reporting + recovery | 4 B | +| `p2_printf_err` | device printf (5 deterministic lines) + passing device `assert` + invalid grid/block launch error reporting + recovery | 4 B | | `p2_bench` | 8192 threads x 250k correctly-rounded FMAs; fixed verifiable value; TIMING ms/GFLOP/s | 32 KB | | `p2_mixed` | mail + vadd + int reduction + 2D grid + 2 streams in one process | peak 96 KB | | `p2_managed_probe` | **expected-failure probe**: cudaMallocManaged outcome recorded, then normal alloc/launch must still work | 4 KB (+4 KB) | @@ -48,7 +51,7 @@ All arithmetic is exact: integer kernels, fp32 add/mul of small integers ## Usage ```bash -./build.sh # build all 10 programs (sm_75, static cudart) +./build.sh # build all 11 programs (sm_75, static cudart) ./oracle.sh # host references -> expected/ # GPU mode if nvidia-smi works, else @@ -87,7 +90,8 @@ The emulated GR path relocates guest data pointers into bounce windows (100 pages total, 112-page win0 shadow, 32 windows, ~33 pages per pointer observed working). This suite stays well inside those caps: largest single pointer 32 KB (8 pages), largest simultaneous footprint 96 KB (24 pages, -`p2_mixed` vadd phase), total live allocations per program <= 4. No managed +`p2_mixed` vadd phase), `p2_regs` 4 KB (1 page), `p2_streams` phase 2 +8 x 4 KB (8 pages). Total live allocations per program <= 4. No managed memory is used except in `p2_managed_probe`, which is informational. ## Special cases @@ -97,6 +101,11 @@ memory is used except in `p2_managed_probe`, which is informational. `INFO` and does not diff its content, but requires it to finish with `STATUS OK` (clean managed-alloc failure, normal alloc still working). A hang is caught by the timeout and fails the suite. +* `p2_printf_err` includes a device `assert` on its passing path only. A + fired (failing) device assert traps the kernel and leaves the CUDA context + poisoned, so it is deliberately not exercised here; the launch-error path + (invalid grid/block) covers error reporting and recovery without killing + the context. * `p2_printf_err` can legitimately be the most fragile test if the emulated path does not support device printf buffering; a timeout/`FAIL` there is a finding, not a harness bug. diff --git a/tools/cuda-ab/p2/expected/SOURCE.txt b/tools/cuda-ab/p2/expected/SOURCE.txt index a138e80..5b08980 100644 --- a/tools/cuda-ab/p2/expected/SOURCE.txt +++ b/tools/cuda-ab/p2/expected/SOURCE.txt @@ -1,5 +1,9 @@ mode=emit -date=2026-10-08T07:01:11Z +date=2026-10-08T15:34:32Z host=majik-envy deterministic CPU computation via --emit-expected (no usable host GPU; rerun ./oracle.sh when one is available) +note=2026-10-08: p2_regs added; p2_streams extended with the independent +4-stream phase; p2_printf_err extended with a passing device assert. All +eleven references regenerated in emit mode; the eight unchanged programs +reproduced their previous references byte-identically. diff --git a/tools/cuda-ab/p2/expected/p2_printf_err.full.txt b/tools/cuda-ab/p2/expected/p2_printf_err.full.txt index e2bf8a4..81995e4 100644 --- a/tools/cuda-ab/p2/expected/p2_printf_err.full.txt +++ b/tools/cuda-ab/p2/expected/p2_printf_err.full.txt @@ -7,4 +7,5 @@ ERR launch grid0 rc=9 name=invalid configuration argument ERR launch block0 rejected=1 ERR cleared rc=0 RESULT errmode grid0_bad=0 block0_bad=0 clear_bad=0 recovered=1 +RESULT assert n=4 bad=0 STATUS OK diff --git a/tools/cuda-ab/p2/expected/p2_printf_err.txt b/tools/cuda-ab/p2/expected/p2_printf_err.txt index e2bf8a4..81995e4 100644 --- a/tools/cuda-ab/p2/expected/p2_printf_err.txt +++ b/tools/cuda-ab/p2/expected/p2_printf_err.txt @@ -7,4 +7,5 @@ ERR launch grid0 rc=9 name=invalid configuration argument ERR launch block0 rejected=1 ERR cleared rc=0 RESULT errmode grid0_bad=0 block0_bad=0 clear_bad=0 recovered=1 +RESULT assert n=4 bad=0 STATUS OK diff --git a/tools/cuda-ab/p2/expected/p2_regs.full.txt b/tools/cuda-ab/p2/expected/p2_regs.full.txt new file mode 100644 index 0000000..c96dbf2 --- /dev/null +++ b/tools/cuda-ab/p2/expected/p2_regs.full.txt @@ -0,0 +1,5 @@ +RESULT regs kernel=lo n=1024 bad=0 +RESULT regs kernel=hi n=1024 bad=0 +RESULT regs kernel=bounds n=1024 bad=0 +RESULT regs kernels=3 bad=0 +STATUS OK diff --git a/tools/cuda-ab/p2/expected/p2_regs.txt b/tools/cuda-ab/p2/expected/p2_regs.txt new file mode 100644 index 0000000..c96dbf2 --- /dev/null +++ b/tools/cuda-ab/p2/expected/p2_regs.txt @@ -0,0 +1,5 @@ +RESULT regs kernel=lo n=1024 bad=0 +RESULT regs kernel=hi n=1024 bad=0 +RESULT regs kernel=bounds n=1024 bad=0 +RESULT regs kernels=3 bad=0 +STATUS OK diff --git a/tools/cuda-ab/p2/expected/p2_streams.full.txt b/tools/cuda-ab/p2/expected/p2_streams.full.txt index 6c60041..30c8a4d 100644 --- a/tools/cuda-ab/p2/expected/p2_streams.full.txt +++ b/tools/cuda-ab/p2/expected/p2_streams.full.txt @@ -1,3 +1,5 @@ RESULT streams nstreams=3 launches=6 n=1024 bad=0 RESULT streams events=3 bad=0 +RESULT streams indep nstreams=4 launches=8 n=1024 bad=0 +RESULT streams indep_events=4 bad=0 STATUS OK diff --git a/tools/cuda-ab/p2/expected/p2_streams.txt b/tools/cuda-ab/p2/expected/p2_streams.txt index 6c60041..30c8a4d 100644 --- a/tools/cuda-ab/p2/expected/p2_streams.txt +++ b/tools/cuda-ab/p2/expected/p2_streams.txt @@ -1,3 +1,5 @@ RESULT streams nstreams=3 launches=6 n=1024 bad=0 RESULT streams events=3 bad=0 +RESULT streams indep nstreams=4 launches=8 n=1024 bad=0 +RESULT streams indep_events=4 bad=0 STATUS OK diff --git a/tools/cuda-ab/p2/run_guest.sh b/tools/cuda-ab/p2/run_guest.sh index 0290df1..50a2b2f 100755 --- a/tools/cuda-ab/p2/run_guest.sh +++ b/tools/cuda-ab/p2/run_guest.sh @@ -45,10 +45,15 @@ if [ ! -x "$VMEXEC" ] && [ ! -f "$VMEXEC" ]; then exit 2 fi -# Build any missing binaries first (bin/ may be absent in a fresh checkout). +# Build any missing or stale binaries first (bin/ may be absent in a fresh +# checkout; sources may be newer than a previously built binary). need_build=0 for name in "${PROGS[@]}"; do - [ -x "$HERE/bin/$name" ] || need_build=1 + bin="$HERE/bin/$name" + if [ ! -x "$bin" ] || [ "$HERE/src/$name.cu" -nt "$bin" ] || + [ "$HERE/src/common.h" -nt "$bin" ]; then + need_build=1 + fi done if [ "$need_build" -eq 1 ]; then "$HERE/build.sh" || exit 2 diff --git a/tools/cuda-ab/p2/src/p2_printf_err.cu b/tools/cuda-ab/p2/src/p2_printf_err.cu index d7f449e..dea2fd5 100644 --- a/tools/cuda-ab/p2/src/p2_printf_err.cu +++ b/tools/cuda-ab/p2/src/p2_printf_err.cu @@ -1,18 +1,23 @@ // SPDX-License-Identifier: BSD-3-Clause /* - * p2_printf_err - device printf plus an intentional launch error path. + * p2_printf_err - device printf, a passing device assert, and an intentional + * launch error path. * * 1. A 4-thread kernel prints deterministic P2PRINT lines from the device. - * 2. A kernel launched with gridDim.x == 0 must be rejected: + * 2. A kernel with assert(n == 4) exercises the device-side assert + * machinery on its passing path (a fired assert is deliberately not + * tested: it traps the kernel and poisons the context). + * 3. A kernel launched with gridDim.x == 0 must be rejected: * cudaErrorInvalidConfiguration (9), and the error must be cleared by * cudaGetLastError. A blockDim.x == 0 launch is exercised too. - * 3. A normal launch after the errors proves error recovery. + * 4. A normal launch after the errors proves error recovery. * * Expected error strings come from the statically linked CUDA runtime, so * they must be identical on host and guest. * Device memory: out 4 B. */ #include "common.h" +#include <cassert> __global__ void hello(int nthreads) { @@ -33,6 +38,13 @@ __global__ void simple(int *out) out[0] = 42; } +__global__ void assertk(int *out, int n) +{ + assert(n == 4); + if (threadIdx.x == 0) + out[0] = n; +} + int main(int argc, char **argv) { if (p2_emit_mode(argc, argv)) { @@ -47,6 +59,7 @@ int main(int argc, char **argv) printf("ERR launch block0 rejected=1\n"); printf("ERR cleared rc=0\n"); printf("RESULT errmode grid0_bad=0 block0_bad=0 clear_bad=0 recovered=1\n"); + printf("RESULT assert n=4 bad=0\n"); return p2_status(0); } @@ -90,10 +103,25 @@ int main(int argc, char **argv) int block0_bad = (e_block == cudaSuccess) ? 1 : 0; int clear_bad = (cleared != cudaSuccess) ? 1 : 0; int recovered = (got == 42) ? 1 : 0; - int bad = grid0_bad + block0_bad + clear_bad + (recovered ? 0 : 1); + int assert_bad = 0; + + assertk<<<1, 1>>>(d_out, 4); + P2_CHECK(cudaGetLastError()); + P2_CHECK(cudaDeviceSynchronize()); + { + int agot = 0; + + P2_CHECK(cudaMemcpy(&agot, d_out, sizeof(int), + cudaMemcpyDeviceToHost)); + assert_bad = (agot == 4) ? 0 : 1; + } + + int bad = grid0_bad + block0_bad + clear_bad + (recovered ? 0 : 1) + + assert_bad; printf("RESULT errmode grid0_bad=%d block0_bad=%d clear_bad=%d recovered=%d\n", grid0_bad, block0_bad, clear_bad, recovered); + printf("RESULT assert n=4 bad=%d\n", assert_bad); cudaFree(d_out); return p2_status(bad); diff --git a/tools/cuda-ab/p2/src/p2_regs.cu b/tools/cuda-ab/p2/src/p2_regs.cu new file mode 100644 index 0000000..d463295 --- /dev/null +++ b/tools/cuda-ab/p2/src/p2_regs.cu @@ -0,0 +1,179 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * p2_regs - register-usage / launch-bounds variants. + * + * Three kernels cover different register-allocation regimes while staying + * exact (unsigned wraparound is well defined; no floating point): + * + * regs_lo __launch_bounds__(256, 4) - strongly capped register use + * (<= 64 registers/thread on the 65536-register file) + * regs_hi 64 independent accumulators x 128 update steps - high + * register pressure (the compiler may spill; correctness is + * unaffected by where the values live) + * regs_bounds __launch_bounds__(128, 4) - a second occupancy target with a + * different block size + * + * Every kernel writes one unsigned per element of a 1024-element buffer; the + * CPU mirrors the exact integer expressions, so the buffer must match + * element for element. INFO lines report the compiler's register counts and + * the queried occupancy (informational; ignored by the A/B comparison). + * + * Device memory: one 4 KB unsigned buffer (1 page). + */ +#include "common.h" + +#define RG_N 1024 + +__global__ void __launch_bounds__(256, 4) regs_lo(unsigned *out, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + out[i] = ((unsigned)i ^ 0x5a5au) + 3u; +} + +__global__ void regs_hi(unsigned *out, int n, unsigned seed) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + unsigned acc[64]; + unsigned s = 0; + + if (i >= n) + return; +#pragma unroll + for (int k = 0; k < 64; k++) + acc[k] = seed + (unsigned)k * ((unsigned)i + 1u); + for (int t = 0; t < 128; t++) { +#pragma unroll + for (int k = 0; k < 64; k++) + acc[k] = acc[k] * 2654435761u + ((unsigned)t ^ (unsigned)k); + } +#pragma unroll + for (int k = 0; k < 64; k++) + s += acc[k]; + out[i] = s; +} + +__global__ void __launch_bounds__(128, 4) regs_bounds(unsigned *out, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + out[i] = (unsigned)(i * 5) + 7u; +} + +static unsigned cpu_lo(int i) +{ + return ((unsigned)i ^ 0x5a5au) + 3u; +} + +static unsigned cpu_hi(int i, unsigned seed) +{ + unsigned acc[64]; + unsigned s = 0; + + for (int k = 0; k < 64; k++) + acc[k] = seed + (unsigned)k * ((unsigned)i + 1u); + for (int t = 0; t < 128; t++) + for (int k = 0; k < 64; k++) + acc[k] = acc[k] * 2654435761u + ((unsigned)t ^ (unsigned)k); + for (int k = 0; k < 64; k++) + s += acc[k]; + return s; +} + +static unsigned cpu_bounds(int i) +{ + return (unsigned)(i * 5) + 7u; +} + +template <typename T> +static void info_attrs(const char *name, T fn, int block) +{ + cudaFuncAttributes a; + int occ = -1; + + if (cudaFuncGetAttributes(&a, fn) != cudaSuccess) + return; + (void)cudaOccupancyMaxActiveBlocksPerMultiprocessor(&occ, fn, block, 0); + printf("INFO regs %s numregs=%d maxthreads=%d occ_blocks_per_sm=%d\n", + name, a.numRegs, a.maxThreadsPerBlock, occ); +} + +int main(int argc, char **argv) +{ + static const unsigned seed = 0x12345678u; + int bad = 0; + + if (p2_emit_mode(argc, argv)) { + printf("RESULT regs kernel=lo n=%d bad=0\n", RG_N); + printf("RESULT regs kernel=hi n=%d bad=0\n", RG_N); + printf("RESULT regs kernel=bounds n=%d bad=0\n", RG_N); + printf("RESULT regs kernels=3 bad=0\n"); + return p2_status(0); + } + + p2_device_info(); + + unsigned *d_out = NULL; + unsigned *got = (unsigned *)malloc(sizeof(unsigned) * RG_N); + + if (!got) + return 1; + P2_CHECK(cudaMalloc(&d_out, sizeof(unsigned) * RG_N)); + + regs_lo<<<(RG_N + 255) / 256, 256>>>(d_out, RG_N); + P2_CHECK(cudaGetLastError()); + P2_CHECK(cudaDeviceSynchronize()); + P2_CHECK(cudaMemcpy(got, d_out, sizeof(unsigned) * RG_N, + cudaMemcpyDeviceToHost)); + { + int sbad = 0; + + for (int i = 0; i < RG_N; i++) + if (got[i] != cpu_lo(i)) + sbad++; + printf("RESULT regs kernel=lo n=%d bad=%d\n", RG_N, sbad); + bad += sbad; + } + + regs_hi<<<(RG_N + 255) / 256, 256>>>(d_out, RG_N, seed); + P2_CHECK(cudaGetLastError()); + P2_CHECK(cudaDeviceSynchronize()); + P2_CHECK(cudaMemcpy(got, d_out, sizeof(unsigned) * RG_N, + cudaMemcpyDeviceToHost)); + { + int sbad = 0; + + for (int i = 0; i < RG_N; i++) + if (got[i] != cpu_hi(i, seed)) + sbad++; + printf("RESULT regs kernel=hi n=%d bad=%d\n", RG_N, sbad); + bad += sbad; + } + + regs_bounds<<<(RG_N + 127) / 128, 128>>>(d_out, RG_N); + P2_CHECK(cudaGetLastError()); + P2_CHECK(cudaDeviceSynchronize()); + P2_CHECK(cudaMemcpy(got, d_out, sizeof(unsigned) * RG_N, + cudaMemcpyDeviceToHost)); + { + int sbad = 0; + + for (int i = 0; i < RG_N; i++) + if (got[i] != cpu_bounds(i)) + sbad++; + printf("RESULT regs kernel=bounds n=%d bad=%d\n", RG_N, sbad); + bad += sbad; + } + + info_attrs("lo", regs_lo, 256); + info_attrs("hi", regs_hi, 256); + info_attrs("bounds", regs_bounds, 128); + + printf("RESULT regs kernels=3 bad=%d\n", bad); + + cudaFree(d_out); + free(got); + return p2_status(bad); +} 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); } |