diff options
Diffstat (limited to 'tools/cuda-ab/perf/perf_workload.cu')
| -rw-r--r-- | tools/cuda-ab/perf/perf_workload.cu | 415 |
1 files changed, 415 insertions, 0 deletions
diff --git a/tools/cuda-ab/perf/perf_workload.cu b/tools/cuda-ab/perf/perf_workload.cu new file mode 100644 index 0000000..e06454e --- /dev/null +++ b/tools/cuda-ab/perf/perf_workload.cu @@ -0,0 +1,415 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * perf_workload.cu - the Open-IOV near-native performance workload. + * + * One binary, two roles: the bare-metal oracle (run on the host with the + * distribution driver loaded) and the measured side (run in the guest on the + * Open-IOV emulated vGPU). Every metric is a stable, machine-readable PERF + * line; every result is verified exactly on both sides, so a perf run can + * never silently measure corrupted work. + * + * Stable output (greppable): + * PERF device=<name> cc=<maj>.<min> sm=<count> + * PERF run_cfg n_warm=.. n_malloc=.. n_thru=.. sizes=.. + * PERF init_ms=<f> + * PERF malloc_first_ms=<f> + * PERF malloc_ms n=.. mean=.. p50=.. p95=.. + * PERF copy dir=htod bytes=.. ms=.. gbps=.. + * PERF copy dir=dtoh bytes=.. ms=.. gbps=.. + * PERF kernel_fma iters=.. ms=.. gflops=.. bad=.. + * PERF launch_first_ms=<f> + * PERF launch_warm n=.. mean=.. p50=.. p95=.. + * PERF throughput n=.. wall_ms=.. kps=.. + * PERF clock64 available=0|1 min=.. median=.. + * PERF checksum fma=.. pattern=.. copy_bad=.. bad=.. + * PERF STATUS OK|FAIL + * + * Build with a static CUDA runtime so the guest needs only the driver's + * libcuda.so, exactly like ab.sh/victory_mail: + * + * nvcc -O2 -arch=sm_75 --cudart=static -o perf_workload perf_workload.cu + * + * Run (either side): + * ./perf_workload [--n-warm=N] [--n-thru=N] [--no-clock64] + */ + +#include <cstdio> +#include <cstdlib> +#include <cstring> +#include <cmath> +#include <ctime> +#include <cuda_runtime.h> + +#define FMA_THREADS 8192 +#define FMA_ITERS 250000 +#define FMA_A 0.9999f +#define FMA_B 1.0e-6f + +#define PAT_THREADS 32768 +#define PAT_SEED 0x9e3779b9u + +#define MALLOC_BYTES (256u * 1024u) +#define N_MALLOC 50 +#define N_COPY 3 +#define CLK_BLOCKS 48 +#define CLK_ITERS 10000 + +static double now_ms(void) +{ + struct timespec ts; + + clock_gettime(CLOCK_MONOTONIC, &ts); + return (double)ts.tv_sec * 1e3 + (double)ts.tv_nsec / 1e6; +} + +static int cmp_double(const void *a, const void *b) +{ + double x = *(const double *)a; + double y = *(const double *)b; + + return (x > y) - (x < y); +} + +static double pct(double *v, int n, double p) +{ + qsort(v, n, sizeof(double), cmp_double); + return v[(int)(p / 100.0 * (n - 1) + 0.5)]; +} + +#define CK(x) do { \ + cudaError_t e_ = (x); \ + if (e_ != cudaSuccess) { \ + printf("PERF ERROR %s at line %d\n", \ + cudaGetErrorString(e_), __LINE__); \ + printf("PERF STATUS FAIL\n"); \ + return 1; \ + } \ +} while (0) + +__global__ void fma_chain(float *out, int iters, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + float x = 0.0f; + + if (i >= n) + return; + for (int t = 0; t < iters; t++) + x = __fmaf_rn(x, FMA_A, FMA_B); + out[i] = x; +} + +static float host_chain(int iters) +{ + float x = 0.0f; + + for (int t = 0; t < iters; t++) + x = fmaf(x, FMA_A, FMA_B); + return x; +} + +__global__ void pattern_kernel(unsigned *out, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i >= n) + return; + out[i] = (unsigned)i * 2654435761u ^ PAT_SEED; +} + +static unsigned host_pattern(int i) +{ + return (unsigned)i * 2654435761u ^ PAT_SEED; +} + +__global__ void clock_probe(unsigned long long *out) +{ + if (threadIdx.x == 0) { + unsigned long long s = clock64(); + float x = 0.0f; + + for (int t = 0; t < CLK_ITERS; t++) + x = __fmaf_rn(x, FMA_A, FMA_B); + out[blockIdx.x * 2 + 0] = s; + out[blockIdx.x * 2 + 1] = clock64(); + } +} + +struct opts { + int n_warm; + int n_thru; + int clock64; +}; + +static int parse_args(int argc, char **argv, struct opts *o) +{ + o->n_warm = 50; + o->n_thru = 100; + o->clock64 = 1; + for (int i = 1; i < argc; i++) { + if (!strncmp(argv[i], "--n-warm=", 9)) + o->n_warm = atoi(argv[i] + 9); + else if (!strncmp(argv[i], "--n-thru=", 9)) + o->n_thru = atoi(argv[i] + 9); + else if (!strcmp(argv[i], "--no-clock64")) + o->clock64 = 0; + else { + printf("usage: %s [--n-warm=N] [--n-thru=N] [--no-clock64]\n", + argv[0]); + return -1; + } + } + if (o->n_warm < 1) + o->n_warm = 1; + if (o->n_thru < 1) + o->n_thru = 1; + return 0; +} + +int main(int argc, char **argv) +{ + struct opts o; + double t_start = now_ms(); + double t0; + float ms; + double samples[256]; + int bad = 0, pat_bad = 0, copy_bad = 0; + float fma_ref = host_chain(FMA_ITERS); + + if (parse_args(argc, argv, &o) != 0) + return 2; + + /* P1: process start -> device ready (the init path). */ + int dev = 0; + struct cudaDeviceProp prop; + + CK(cudaGetDeviceCount(&dev)); + if (dev < 1) { + printf("PERF ERROR no CUDA devices\n"); + printf("PERF STATUS FAIL\n"); + return 1; + } + CK(cudaSetDevice(0)); + CK(cudaGetDeviceProperties(&prop, 0)); + printf("PERF device=%s cc=%d.%d sm=%d\n", prop.name, + prop.major, prop.minor, prop.multiProcessorCount); + + /* Touch the device so init_ms covers the full lazy-init path. */ + CK(cudaFree(0)); + printf("PERF init_ms=%.3f\n", now_ms() - t_start); + printf("PERF run_cfg n_warm=%d n_malloc=%d n_thru=%d sizes=4096,262144,4194304,16777216\n", + o.n_warm, N_MALLOC, o.n_thru); + + /* P2: cudaMalloc first-call and steady-state latency. */ + void *d_m = NULL; + + t0 = now_ms(); + CK(cudaMalloc(&d_m, MALLOC_BYTES)); + printf("PERF malloc_first_ms=%.3f\n", now_ms() - t0); + CK(cudaFree(d_m)); + + for (int i = 0; i < 8; i++) { + CK(cudaMalloc(&d_m, MALLOC_BYTES)); + CK(cudaFree(d_m)); + } + for (int i = 0; i < N_MALLOC; i++) { + t0 = now_ms(); + CK(cudaMalloc(&d_m, MALLOC_BYTES)); + CK(cudaFree(d_m)); + samples[i] = now_ms() - t0; + } + { + double s2[256]; + double a = 0; + + memcpy(s2, samples, sizeof(double) * N_MALLOC); + for (int i = 0; i < N_MALLOC; i++) + a += samples[i]; + printf("PERF malloc_ms n=%d mean=%.3f p50=%.3f p95=%.3f\n", + N_MALLOC, a / N_MALLOC, + pct(s2, N_MALLOC, 50), pct(s2, N_MALLOC, 95)); + } + + /* Verification pattern kernel + first launch. */ + unsigned *d_pat = NULL, *h_pat; + + h_pat = (unsigned *)malloc(sizeof(unsigned) * PAT_THREADS); + CK(cudaMalloc(&d_pat, sizeof(unsigned) * PAT_THREADS)); + + cudaEvent_t e0, e1; + + CK(cudaEventCreate(&e0)); + CK(cudaEventCreate(&e1)); + CK(cudaEventRecord(e0)); + pattern_kernel<<<(PAT_THREADS + 255) / 256, 256>>>(d_pat, PAT_THREADS); + CK(cudaGetLastError()); + CK(cudaEventRecord(e1)); + CK(cudaDeviceSynchronize()); + CK(cudaEventElapsedTime(&ms, e0, e1)); + printf("PERF launch_first_ms=%.3f\n", (double)ms); + CK(cudaMemcpy(h_pat, d_pat, sizeof(unsigned) * PAT_THREADS, + cudaMemcpyDeviceToHost)); + for (int i = 0; i < PAT_THREADS; i++) + if (h_pat[i] != host_pattern(i)) + pat_bad++; + + /* P4: warm launch latency (enqueue -> sync completion per launch). */ + if (o.n_warm > 256) + o.n_warm = 256; + for (int i = 0; i < o.n_warm; i++) { + CK(cudaEventRecord(e0)); + pattern_kernel<<<(PAT_THREADS + 255) / 256, 256>>>(d_pat, + PAT_THREADS); + CK(cudaGetLastError()); + CK(cudaEventRecord(e1)); + CK(cudaDeviceSynchronize()); + CK(cudaEventElapsedTime(&ms, e0, e1)); + samples[i] = (double)ms; + } + { + double s2[256]; + double a = 0; + + memcpy(s2, samples, sizeof(double) * o.n_warm); + for (int i = 0; i < o.n_warm; i++) + a += samples[i]; + printf("PERF launch_warm n=%d mean=%.3f p50=%.3f p95=%.3f\n", + o.n_warm, a / o.n_warm, + pct(s2, o.n_warm, 50), pct(s2, o.n_warm, 95)); + } + + /* P5/P9: compute kernel vs the bare-metal oracle (bit-exact). */ + float *d_fma = NULL, *h_fma; + + h_fma = (float *)malloc(sizeof(float) * FMA_THREADS); + CK(cudaMalloc(&d_fma, sizeof(float) * FMA_THREADS)); + CK(cudaEventRecord(e0)); + fma_chain<<<FMA_THREADS / 256, 256>>>(d_fma, FMA_ITERS, FMA_THREADS); + CK(cudaGetLastError()); + CK(cudaEventRecord(e1)); + CK(cudaDeviceSynchronize()); + CK(cudaEventElapsedTime(&ms, e0, e1)); + printf("PERF kernel_fma iters=%d ms=%.3f gflops=%.2f\n", + FMA_ITERS, (double)ms, + 2.0 * (double)FMA_THREADS * (double)FMA_ITERS / + ((double)ms * 1.0e6)); + CK(cudaMemcpy(h_fma, d_fma, sizeof(float) * FMA_THREADS, + cudaMemcpyDeviceToHost)); + for (int i = 0; i < FMA_THREADS; i++) + if (h_fma[i] != fma_ref) + bad++; + + /* P10: steady-state throughput (enqueue burst, one sync). */ + CK(cudaEventRecord(e0)); + for (int i = 0; i < o.n_thru; i++) { + pattern_kernel<<<(PAT_THREADS + 255) / 256, 256>>>(d_pat, + PAT_THREADS); + CK(cudaGetLastError()); + } + CK(cudaEventRecord(e1)); + CK(cudaDeviceSynchronize()); + CK(cudaEventElapsedTime(&ms, e0, e1)); + printf("PERF throughput n=%d wall_ms=%.3f kps=%.2f\n", + o.n_thru, (double)ms, + (double)o.n_thru / ((double)ms / 1.0e3)); + + /* P7/P8: HtoD and DtoH copy bandwidth, and a small round-trip check. */ + { + const size_t sizes[] = { 4096u, 262144u, 4194304u, 16777216u }; + unsigned char *h_a, *h_b, *d_a; + + h_a = (unsigned char *)malloc(sizes[3]); + h_b = (unsigned char *)malloc(sizes[3]); + if (!h_a || !h_b) { + printf("PERF ERROR host alloc failed\n"); + printf("PERF STATUS FAIL\n"); + return 1; + } + CK(cudaMalloc((void **)&d_a, sizes[3])); + for (int s = 0; s < 4; s++) { + size_t n = sizes[s]; + double best_h = 1e30, best_d = 1e30; + + for (int i = 0; i < (int)n; i++) + h_a[i] = (unsigned char)(i * 131 + s); + for (int rep = 0; rep <= N_COPY; rep++) { + CK(cudaEventRecord(e0)); + CK(cudaMemcpy(d_a, h_a, n, cudaMemcpyHostToDevice)); + CK(cudaEventRecord(e1)); + CK(cudaDeviceSynchronize()); + CK(cudaEventElapsedTime(&ms, e0, e1)); + if (rep > 0 && (double)ms < best_h) + best_h = (double)ms; + } + for (int rep = 0; rep <= N_COPY; rep++) { + CK(cudaEventRecord(e0)); + CK(cudaMemcpy(h_b, d_a, n, cudaMemcpyDeviceToHost)); + CK(cudaEventRecord(e1)); + CK(cudaDeviceSynchronize()); + CK(cudaEventElapsedTime(&ms, e0, e1)); + if (rep > 0 && (double)ms < best_d) + best_d = (double)ms; + } + printf("PERF copy dir=htod bytes=%zu ms=%.4f gbps=%.3f\n", + n, best_h, (double)n / (best_h * 1.0e6)); + printf("PERF copy dir=dtoh bytes=%zu ms=%.4f gbps=%.3f\n", + n, best_d, (double)n / (best_d * 1.0e6)); + if (s == 0 && memcmp(h_a, h_b, n) != 0) + copy_bad++; + } + free(h_a); + free(h_b); + cudaFree(d_a); + } + + /* Optional clock64 cross-check (SM cycle counter, not wall time). */ + if (o.clock64) { + unsigned long long *d_clk, *h_clk; + + h_clk = (unsigned long long *)malloc( + sizeof(unsigned long long) * CLK_BLOCKS * 2); + CK(cudaMalloc((void **)&d_clk, + sizeof(unsigned long long) * CLK_BLOCKS * 2)); + CK(cudaMemset(d_clk, 0, + sizeof(unsigned long long) * CLK_BLOCKS * 2)); + clock_probe<<<CLK_BLOCKS, 32>>>(d_clk); + CK(cudaGetLastError()); + CK(cudaDeviceSynchronize()); + CK(cudaMemcpy(h_clk, d_clk, + sizeof(unsigned long long) * CLK_BLOCKS * 2, + cudaMemcpyDeviceToHost)); + { + double cyc[CLK_BLOCKS]; + int n = 0; + + for (int i = 0; i < CLK_BLOCKS; i++) { + if (h_clk[i * 2 + 1] > h_clk[i * 2] && + h_clk[i * 2] != 0) + cyc[n++] = (double)(h_clk[i * 2 + 1] - + h_clk[i * 2]); + } + if (n > 0) + printf("PERF clock64 available=1 min=%.0f median=%.0f\n", + pct(cyc, n, 0), pct(cyc, n, 50)); + else + printf("PERF clock64 available=0 min=0 median=0\n"); + } + free(h_clk); + cudaFree(d_clk); + } else { + printf("PERF clock64 available=0 min=0 median=0\n"); + } + + printf("PERF checksum fma=%.9g pattern=%u copy_bad=%d bad=%d\n", + (double)fma_ref, host_pattern(PAT_THREADS - 1), copy_bad, + bad + pat_bad); + printf("PERF STATUS %s\n", + (bad + pat_bad + copy_bad) == 0 ? "OK" : "FAIL"); + + cudaEventDestroy(e0); + cudaEventDestroy(e1); + cudaFree(d_pat); + cudaFree(d_fma); + free(h_pat); + free(h_fma); + return (bad + pat_bad + copy_bad) == 0 ? 0 : 1; +} |