diff options
Diffstat (limited to 'tools/cuda-ab')
| -rwxr-xr-x | tools/cuda-ab/ab.sh | 1 | ||||
| -rw-r--r-- | tools/cuda-ab/smiley.cu | 230 |
2 files changed, 231 insertions, 0 deletions
diff --git a/tools/cuda-ab/ab.sh b/tools/cuda-ab/ab.sh index 8ae05ee..5a290e4 100755 --- a/tools/cuda-ab/ab.sh +++ b/tools/cuda-ab/ab.sh @@ -86,6 +86,7 @@ host) run_host ;; guest) run_guest ;; compare) compare ;; all) run_host; run_guest; compare ;; +push) push_bin ;; victory) WORKLOAD=victory_mail SRC="$HERE/$WORKLOAD.cu" diff --git a/tools/cuda-ab/smiley.cu b/tools/cuda-ab/smiley.cu new file mode 100644 index 0000000..42770cc --- /dev/null +++ b/tools/cuda-ab/smiley.cu @@ -0,0 +1,230 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * smiley.cu - vectorAdd draws a smiley face. + * + * Four vectorAdd passes execute on the GPU; each computes one scalar field + * + * c[i] = a[i] + b[i] + * + * over the same n = 32768 elements laid out as a 256x128 sample grid: + * + * pass 0 a=u^2 b=v^2 -> face disk field + * pass 1 a=(u+xe)^2 b=(v+ye)^2 -> left-eye disk field + * pass 2 a=(u-xe)^2 b=(v+ye)^2 -> right-eye disk field + * pass 3 a=u^2 b=(v-ym)^2 -> mouth annulus field + * + * with u,v centre-relative, aspect-corrected coordinates. The host copies + * the four fields back and thresholds them to render a 64x32 ASCII smiley + * (4x4 samples per character cell). The artwork therefore comes out of GPU + * addition results only - no GPU-side drawing, no CPU-side math on pixels + * beyond the thresholds. + * + * Output lines are stable and greppable (matches tools/cuda-ab/ab.sh): + * + * DEVICE <name> cc=<major>.<minor> sm=<count> mem=<bytes> + * MALLOC bytes=<n> rc=<0|err> + * VADD passes=4 n=<count> bad=<mismatches> checksum=<sum> + * RESULT cells=<W>x<H> inked=<cells> checksum=<sum> + * 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 smiley smiley.cu + * + * Then, from the guest: /tmp/smiley + */ + +#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]; +} + +/* sample grid */ +#define GW 256 +#define GH 128 +#define GN (GW * GH) + +/* character cell = 4x4 samples */ +#define CW 64 +#define CH 32 +#define CELL 4 + +/* feature geometry, in aspect-corrected grid units (u ~ +/-127, v ~ +/-127) */ +#define FACE_R2 10000.0f /* face disk: u^2+v^2 <= 100^2 */ +#define EYE_X 36.0f +#define EYE_Y 30.0f +#define EYE_R2 144.0f /* eyes: r=12 */ +#define MOUTH_Y 24.0f +#define MOUTH_IN 2025.0f /* 45^2 */ +#define MOUTH_OUT 3025.0f /* 55^2 */ +#define MOUTH_V 40.0f + +int main(void) +{ + cudaDeviceProp prop; + float *h_a, *h_b, *h_out, *fields; + float *d_a, *d_b, *d_c; + size_t bytes = (size_t)GN * sizeof(float); + double checksum = 0.0; + int bad = 0, pass, i; + + 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); + + { + cudaError_t e_a = cudaMalloc(&d_a, bytes); + cudaError_t e_b = e_a == cudaSuccess ? + cudaMalloc(&d_b, bytes) : e_a; + cudaError_t e_c = e_b == cudaSuccess ? + cudaMalloc(&d_c, bytes) : e_b; + + printf("MALLOC bytes=%zu rc=%d\n", bytes, (int)e_a); + if (e_a != cudaSuccess) { + size_t fm = 0, tm = 0; + + cudaMemGetInfo(&fm, &tm); + printf("MEMINFO free=%llu total=%llu\n", + (unsigned long long)fm, + (unsigned long long)tm); + } + CHECK(e_a); + CHECK(e_b); + CHECK(e_c); + } + + h_a = (float *)malloc(bytes); + h_b = (float *)malloc(bytes); + h_out = (float *)malloc(bytes); + fields = (float *)malloc((size_t)4 * bytes); + if (!h_a || !h_b || !h_out || !fields) + return 1; + + for (pass = 0; pass < 4; pass++) { + for (i = 0; i < GN; i++) { + int x = i % GW, y = i / GW; + float u = (float)x - 127.5f; + float v = ((float)y - 63.5f) * 2.0f; + float pa, pb; + + switch (pass) { + case 0: /* face */ + pa = u * u; pb = v * v; + break; + case 1: /* left eye */ + pa = (u + EYE_X) * (u + EYE_X); + pb = (v + EYE_Y) * (v + EYE_Y); + break; + case 2: /* right eye */ + pa = (u - EYE_X) * (u - EYE_X); + pb = (v + EYE_Y) * (v + EYE_Y); + break; + default: /* mouth */ + pa = u * u; + pb = (v - MOUTH_Y) * (v - MOUTH_Y); + break; + } + h_a[i] = pa; + h_b[i] = pb; + } + + CHECK(cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice)); + CHECK(cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice)); + vadd<<<(GN + 255) / 256, 256>>>(d_a, d_b, d_c, GN); + CHECK(cudaGetLastError()); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h_out, d_c, bytes, cudaMemcpyDeviceToHost)); + + for (i = 0; i < GN; i++) { + if (h_out[i] != h_a[i] + h_b[i]) + bad++; + checksum += (double)h_out[i]; + } + for (i = 0; i < GN; i++) + fields[(size_t)pass * GN + i] = h_out[i]; + } + + printf("VADD passes=4 n=%d bad=%d checksum=%.0f\n", GN, bad, checksum); + if (bad != 0) { + printf("STATUS FAIL\n"); + return 1; + } + + { + static char art[CH][CW + 1]; + int inked = 0, cy, cx; + + for (cy = 0; cy < CH; cy++) { + for (cx = 0; cx < CW; cx++) { + int hit = 0, sy, sx; + const float *f0 = fields; + const float *f1 = fields + GN; + const float *f2 = fields + 2 * GN; + const float *f3 = fields + 3 * GN; + + for (sy = 0; sy < CELL; sy++) { + int y = cy * CELL + sy; + float v = ((float)y - 63.5f) * 2.0f; + + for (sx = 0; sx < CELL; sx++) { + int x = cx * CELL + sx; + int idx = y * GW + x; + int inside; + + inside = f0[idx] <= FACE_R2; + if (f1[idx] <= EYE_R2 || + f2[idx] <= EYE_R2) + inside = 0; + if (f3[idx] >= MOUTH_IN && + f3[idx] <= MOUTH_OUT && + v >= MOUTH_V) + inside = 0; + hit += inside; + } + } + art[cy][cx] = hit >= CELL * CELL / 2 ? '#' : + ' '; + inked += art[cy][cx] == '#'; + } + art[cy][CW] = 0; + } + + printf("RESULT cells=%dx%d inked=%d checksum=%.0f\n", + CW, CH, inked, checksum); + for (cy = 0; cy < CH; cy++) + printf("%s\n", art[cy]); + } + + printf("STATUS OK\n"); + CHECK(cudaFree(d_a)); + CHECK(cudaFree(d_b)); + CHECK(cudaFree(d_c)); + free(h_a); + free(h_b); + free(h_out); + free(fields); + return 0; +} |