1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
|
// SPDX-License-Identifier: BSD-3-Clause
/*
* p2_bench - compute-heavy FMA benchmark with a fixed, verifiable result.
*
* 8192 threads each run a 250,000-iteration serial chain of correctly
* rounded FMAs (__fmaf_rn on device, fmaf on host), producing 2.05e9 FMAs.
* Both implementations are IEEE-correctly-rounded, so the per-thread final
* value must match the host exactly (bad=0); the same value is printed for
* the A/B comparison. A TIMING line (ignored by the harness comparison)
* reports wall-clock ms and GFLOP/s for the performance demonstration.
* Device memory: 32 KB output (8 pages).
*/
#include "common.h"
#include <math.h>
#define BENCH_THREADS 8192
#define BENCH_ITERS 250000
#define BENCH_A 0.9999f
#define BENCH_B 1.0e-6f
__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, BENCH_A, BENCH_B);
out[i] = x;
}
static float host_chain(int iters)
{
float x = 0.0f;
for (int t = 0; t < iters; t++)
x = fmaf(x, BENCH_A, BENCH_B);
return x;
}
int main(int argc, char **argv)
{
float ref = host_chain(BENCH_ITERS);
if (p2_emit_mode(argc, argv)) {
printf("RESULT bench threads=%d iters=%d bad=0 value=%.9g\n",
BENCH_THREADS, BENCH_ITERS, (double)ref);
printf("TIMING bench ms=0.000 gflops=0.00 (emit mode: no GPU)\n");
return p2_status(0);
}
p2_device_info();
float *d_out = NULL, *got = (float *)malloc(sizeof(float) * BENCH_THREADS);
if (!got)
return 1;
P2_CHECK(cudaMalloc(&d_out, sizeof(float) * BENCH_THREADS));
cudaEvent_t t0, t1;
float ms = 0.0f;
P2_CHECK(cudaEventCreate(&t0));
P2_CHECK(cudaEventCreate(&t1));
P2_CHECK(cudaEventRecord(t0));
fma_chain<<<BENCH_THREADS / 256, 256>>>(d_out, BENCH_ITERS,
BENCH_THREADS);
P2_CHECK(cudaGetLastError());
P2_CHECK(cudaEventRecord(t1));
P2_CHECK(cudaDeviceSynchronize());
P2_CHECK(cudaEventElapsedTime(&ms, t0, t1));
P2_CHECK(cudaMemcpy(got, d_out, sizeof(float) * BENCH_THREADS,
cudaMemcpyDeviceToHost));
int bad = 0;
for (int i = 0; i < BENCH_THREADS; i++)
if (got[i] != ref)
bad++;
printf("RESULT bench threads=%d iters=%d bad=%d value=%.9g\n",
BENCH_THREADS, BENCH_ITERS, bad, (double)got[0]);
if (ms > 0.0f) {
double flops = 2.0 * (double)BENCH_THREADS * (double)BENCH_ITERS;
printf("TIMING bench ms=%.3f gflops=%.2f\n", (double)ms,
flops / ((double)ms * 1.0e6));
}
cudaEventDestroy(t0);
cudaEventDestroy(t1);
cudaFree(d_out);
free(got);
return p2_status(bad);
}
|
