diff options
Diffstat (limited to 'tools/cuda-ab/p2/src/p2_params.cu')
| -rw-r--r-- | tools/cuda-ab/p2/src/p2_params.cu | 185 |
1 files changed, 185 insertions, 0 deletions
diff --git a/tools/cuda-ab/p2/src/p2_params.cu b/tools/cuda-ab/p2/src/p2_params.cu new file mode 100644 index 0000000..5dd942c --- /dev/null +++ b/tools/cuda-ab/p2/src/p2_params.cu @@ -0,0 +1,185 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * p2_params - kernel parameter-count and parameter-type coverage. + * + * k0..k12 have 0..12 value parameters respectively (plus the output + * pointer), mixing int, float, double and device-pointer types. Every + * kernel writes one int slot of a 13-int device buffer; the CPU reference + * mirrors each expression exactly and the buffer must match slot for slot. + * + * Device memory: out 52 B, p 8 B, pf 4 B (1 page). + */ +#include "common.h" + +#define NK 13 + +__global__ void k0(int *out) +{ + out[0] = 7; +} + +__global__ void k1(int *out, int a) +{ + out[1] = a + 1; +} + +__global__ void k2(int *out, int a, int b) +{ + out[2] = a + b; +} + +__global__ void k3(int *out, int a, int b, float f1) +{ + out[3] = a + b + (int)f1; +} + +__global__ void k4(int *out, int a, int b, float f1, double c1) +{ + out[4] = a + b + (int)f1 + (int)c1; +} + +__global__ void k5(int *out, int a, int b, float f1, double c1, int d) +{ + out[5] = a + b + (int)f1 + (int)c1 + d; +} + +__global__ void k6(int *out, int a, int b, float f1, double c1, int d, + const int *p) +{ + out[6] = a + b + (int)f1 + (int)c1 + d + p[0]; +} + +__global__ void k7(int *out, int a, int b, float f1, double c1, int d, + const int *p, const int *q) +{ + out[7] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0]; +} + +__global__ void k8(int *out, int a, int b, float f1, double c1, int d, + const int *p, const int *q, int n) +{ + out[8] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0] + n; +} + +__global__ void k9(int *out, int a, int b, float f1, double c1, int d, + const int *p, const int *q, int n, float f2) +{ + out[9] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0] + n + (int)f2; +} + +__global__ void k10(int *out, int a, int b, float f1, double c1, int d, + const int *p, const int *q, int n, float f2, + double c2) +{ + out[10] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0] + n + + (int)f2 + (int)c2; +} + +__global__ void k11(int *out, int a, int b, float f1, double c1, int d, + const int *p, const int *q, int n, float f2, + double c2, float s) +{ + out[11] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0] + n + + (int)f2 + (int)c2 + (int)s; +} + +__global__ void k12(int *out, int a, int b, float f1, double c1, int d, + const int *p, const int *q, int n, float f2, + double c2, float s, int e) +{ + out[12] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0] + n + + (int)f2 + (int)c2 + (int)s + e; +} + +/* CPU mirror of the 13 kernel expressions; keeps expected values exact. */ +static void expected(int *exp, int a, int b, float f1, double c1, int d, + const int *p, const int *q, int n, float f2, double c2, + float s, int e) +{ + exp[0] = 7; + exp[1] = a + 1; + exp[2] = a + b; + exp[3] = a + b + (int)f1; + exp[4] = a + b + (int)f1 + (int)c1; + exp[5] = a + b + (int)f1 + (int)c1 + d; + exp[6] = a + b + (int)f1 + (int)c1 + d + p[0]; + exp[7] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0]; + exp[8] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0] + n; + exp[9] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0] + n + (int)f2; + exp[10] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0] + n + + (int)f2 + (int)c2; + exp[11] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0] + n + + (int)f2 + (int)c2 + (int)s; + exp[12] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0] + n + + (int)f2 + (int)c2 + (int)s + e; +} + +static int checksum(const int *v) +{ + int sum = 0; + + for (int i = 0; i < NK; i++) + sum += v[i]; + return sum; +} + +int main(int argc, char **argv) +{ + int a = 3, b = 5, d = 11, n = 2, e = 4; + float f1 = 6.5f, f2 = 8.5f, s = 2.5f; + double c1 = 2.5, c2 = 3.5; + int hp[2] = { 5, 6 }, hq[2] = { 6, 7 }; + int exp[NK]; + + expected(exp, a, b, f1, c1, d, hp, hq, n, f2, c2, s, e); + + if (p2_emit_mode(argc, argv)) { + printf("RESULT params kernels=%d nparam_max=12 bad=0 checksum=%d\n", + NK, checksum(exp)); + return p2_status(0); + } + + p2_device_info(); + + int *d_out = NULL, *got = (int *)malloc(sizeof(int) * NK); + int *dp = NULL, *dq = NULL; + + if (!got) + return 1; + P2_CHECK(cudaMalloc(&d_out, sizeof(int) * NK)); + P2_CHECK(cudaMalloc(&dp, sizeof(int) * 2)); + P2_CHECK(cudaMalloc(&dq, sizeof(int) * 2)); + P2_CHECK(cudaMemcpy(dp, hp, sizeof(int) * 2, cudaMemcpyHostToDevice)); + P2_CHECK(cudaMemcpy(dq, hq, sizeof(int) * 2, cudaMemcpyHostToDevice)); + P2_CHECK(cudaMemset(d_out, 0, sizeof(int) * NK)); + + k0<<<1, 1>>>(d_out); + k1<<<1, 1>>>(d_out, a); + k2<<<1, 1>>>(d_out, a, b); + k3<<<1, 1>>>(d_out, a, b, f1); + k4<<<1, 1>>>(d_out, a, b, f1, c1); + k5<<<1, 1>>>(d_out, a, b, f1, c1, d); + k6<<<1, 1>>>(d_out, a, b, f1, c1, d, dp); + k7<<<1, 1>>>(d_out, a, b, f1, c1, d, dp, dq); + k8<<<1, 1>>>(d_out, a, b, f1, c1, d, dp, dq, n); + k9<<<1, 1>>>(d_out, a, b, f1, c1, d, dp, dq, n, f2); + k10<<<1, 1>>>(d_out, a, b, f1, c1, d, dp, dq, n, f2, c2); + k11<<<1, 1>>>(d_out, a, b, f1, c1, d, dp, dq, n, f2, c2, s); + k12<<<1, 1>>>(d_out, a, b, f1, c1, d, dp, dq, n, f2, c2, s, e); + P2_CHECK(cudaGetLastError()); + P2_CHECK(cudaDeviceSynchronize()); + P2_CHECK(cudaMemcpy(got, d_out, sizeof(int) * NK, + cudaMemcpyDeviceToHost)); + + int bad = 0; + + for (int i = 0; i < NK; i++) + if (got[i] != exp[i]) + bad++; + printf("RESULT params kernels=%d nparam_max=12 bad=%d checksum=%d\n", + NK, bad, checksum(got)); + + cudaFree(d_out); cudaFree(dp); cudaFree(dq); + free(got); + return p2_status(bad); +} |