diff options
| -rwxr-xr-x | tools/cuda-ab/run/wedge_ab | bin | 0 -> 825144 bytes | |||
| -rw-r--r-- | tools/cuda-ab/wedge_ab.cu | 100 | ||||
| -rw-r--r-- | tools/lg/frame_reader.py | 5 |
3 files changed, 103 insertions, 2 deletions
diff --git a/tools/cuda-ab/run/wedge_ab b/tools/cuda-ab/run/wedge_ab Binary files differnew file mode 100755 index 0000000..23a90aa --- /dev/null +++ b/tools/cuda-ab/run/wedge_ab diff --git a/tools/cuda-ab/wedge_ab.cu b/tools/cuda-ab/wedge_ab.cu new file mode 100644 index 0000000..2d88f92 --- /dev/null +++ b/tools/cuda-ab/wedge_ab.cu @@ -0,0 +1,100 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * wedge_ab.cu - discriminate the wedge trigger: is it the second launch of + * the SAME program, or the second launch of the context regardless of + * program? + * + * Sequence (each launch verified against mutated inputs): + * 1. vadd (program A) + * 2. vsub (program B, different code) + * 3. vadd (program A again) + * 4. vsub (program B again) + * + * If launch 2 (B) executes -> the trigger is same-program reuse. + * If launch 2 (B) wedges -> the trigger is the second launch in the context. + * + * Build: nvcc -O2 --cudart=static -o wedge_ab wedge_ab.cu + */ + +#include <cstdio> +#include <cstdlib> +#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 vsub(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]; +} + +int main(void) +{ + const int n = 1 << 15; + size_t bytes = (size_t)n * sizeof(float); + float *ha, *hb, *hc; + float *da, *db, *dc; + int i, bad, step; + + setvbuf(stdout, NULL, _IONBF, 0); + ha = (float *)malloc(bytes); + hb = (float *)malloc(bytes); + hc = (float *)malloc(bytes); + if (!ha || !hb || !hc) + return 1; + for (i = 0; i < n; i++) { + ha[i] = (float)(i % 977) * 0.25f; + hb[i] = (float)(i % 251) * 0.5f; + } + CHECK(cudaMalloc(&da, bytes)); + CHECK(cudaMalloc(&db, bytes)); + CHECK(cudaMalloc(&dc, bytes)); + + for (step = 0; step < 4; step++) { + const char *which = (step & 1) ? "vsub" : "vadd"; + int is_sub = step & 1; + + if (step > 0) { + ha[step] += 1.0f; + hb[step * 7 + 3] += 2.0f; + } + CHECK(cudaMemcpy(da, ha, bytes, cudaMemcpyHostToDevice)); + CHECK(cudaMemcpy(db, hb, bytes, cudaMemcpyHostToDevice)); + if (is_sub) + vsub<<<(n + 255) / 256, 256>>>(da, db, dc, n); + else + vadd<<<(n + 255) / 256, 256>>>(da, db, dc, n); + CHECK(cudaGetLastError()); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(hc, dc, bytes, cudaMemcpyDeviceToHost)); + bad = 0; + for (i = 0; i < n; i++) { + float expect = is_sub ? ha[i] - hb[i] : ha[i] + hb[i]; + + if (hc[i] != expect) + bad++; + } + printf("STEP %d %s bad=%d/%d\n", step, which, bad, n); + } + printf("STATUS OK\n"); + return 0; +} diff --git a/tools/lg/frame_reader.py b/tools/lg/frame_reader.py index 5598667..f8bab40 100644 --- a/tools/lg/frame_reader.py +++ b/tools/lg/frame_reader.py @@ -30,11 +30,12 @@ def main(): if not args.shm: import glob - cands = sorted(glob.glob("/dev/shm/kvmfr-*")) + cands = [p for p in glob.glob("/dev/shm/kvmfr-*") + if os.path.getsize(p) > 0] if not cands: print("NO_SHM_FILE", file=sys.stderr) return 1 - args.shm = cands[0] + args.shm = max(cands, key=os.path.getsize) os.makedirs(args.out, exist_ok=True) print("SHM %s" % args.shm) |