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
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
|
// SPDX-License-Identifier: BSD-3-Clause
/*
* r18_anim.cu - GPU-computed animation for R18 (v3: single-page kernel).
*
* Renders NFRAMES frames of an animated plasma with one launch per frame.
* The emulated forward path delivers the kernel through ONE 4 KB code page
* (the window shadow copies the program page; multi-page kernels execute off
* the end of the mapped page - the 904-instruction v2 kernel faulted with
* "Illegal Instruction Encoding" on the real engine). This version keeps the
* SASS text under 4096 bytes: 512x360 RGBA8 (power-of-two width for cheap
* index math), grid-stride loop, three MUFU.SIN under --use_fast_math.
*
* ONE 0.9 MB device frame buffer is reused for every frame (small-alloc
* regime, matching the proven CUDA path); after each launch the frame is
* copied back (DtoH) and APPENDED to /dev/shm with a magic header so the
* host can find it by scanning the guest RAM file. The file is
*
* "R18FRM01" u32 count u32 width u32 height u32 stride
* then count*height*stride bytes of RGBA8
*
* The header is written first with the final count and bytes are appended
* (flushed) per frame, so a host-side harvest can recover every completed
* frame even if the client or guest wedges mid-run.
*
* Build (host):
* nvcc -O2 --use_fast_math --cudart=static -arch=sm_75 -o r18_anim r18_anim.cu
* Usage: r18_anim [frames] [outpath] [blocks]
*/
#include <cstdio>
#include <cstdlib>
#include <cstdint>
#include <cstring>
#include <ctime>
#include <cuda_runtime.h>
#define WIDTH 512
#define HEIGHT 360
#define NFRAMES 60
#define OUTPATH "/dev/shm/r18frames.bin"
#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)
/* 3 MUFU.SIN only - the text must stay under one 4 KB page. */
__global__ void plasma_kernel(uchar4 *fb, int w, int h, int frame)
{
unsigned stride = gridDim.x * blockDim.x;
unsigned xmask = (unsigned)w - 1u; /* w is a power of two */
int idx;
for (idx = blockIdx.x * blockDim.x + threadIdx.x; idx < w * h;
idx += stride) {
float u = (float)((unsigned)idx & xmask) * (1.0f / 512.0f);
float v = (float)((unsigned)idx >> 9) * (1.0f / 360.0f);
float t = (float)frame * 0.10f;
float p = sinf(u * 9.0f + t * 2.0f)
+ sinf(v * 7.0f - t * 2.0f);
float s = sinf(p * 1.5f);
p *= 0.5f; /* -1 .. 1 */
fb[idx] = make_uchar4((unsigned char)(127.5f + 127.5f * s),
(unsigned char)(127.5f + 127.5f * p),
(unsigned char)(127.5f - 127.5f * s),
255);
}
}
static uint64_t fnv1a(const unsigned char *p, size_t n)
{
uint64_t h = 1469598103934665603ull;
while (n--) {
h ^= *p++;
h *= 1099511628211ull;
}
return h;
}
static double now_ms(void)
{
struct timespec ts;
clock_gettime(CLOCK_MONOTONIC, &ts);
return (double)ts.tv_sec * 1e3 + (double)ts.tv_nsec / 1e6;
}
int main(int argc, char **argv)
{
int nframes = NFRAMES;
const char *outpath = OUTPATH;
size_t fbytes = (size_t)WIDTH * HEIGHT * 4;
uchar4 *d_fb = NULL;
unsigned char *h_fb = NULL;
double km = 0.0, cm = 0.0, t0, t1, wall0 = now_ms();
int blocks = 720;
int f, dev = -1;
FILE *fp;
uint32_t hdr[4];
if (argc > 1)
nframes = atoi(argv[1]);
if (nframes <= 0 || nframes > 60000)
nframes = NFRAMES;
if (argc > 2)
outpath = argv[2];
if (argc > 3)
blocks = atoi(argv[3]);
if (blocks <= 0 || blocks > 4096)
blocks = 720;
setvbuf(stdout, NULL, _IONBF, 0);
CHECK(cudaGetDevice(&dev));
{
cudaDeviceProp prop;
if (cudaGetDeviceProperties(&prop, dev) == cudaSuccess)
printf("DEVICE %s cc=%d.%d\n", prop.name,
prop.major, prop.minor);
}
printf("ANIM start frames=%d %dx%d blocks=%d out=%s\n",
nframes, WIDTH, HEIGHT, blocks, outpath);
CHECK(cudaMalloc(&d_fb, fbytes));
h_fb = (unsigned char *)malloc(fbytes);
if (!h_fb) {
printf("HOST_ALLOC_FAIL\nSTATUS FAIL\n");
return 1;
}
fp = fopen(outpath, "wb");
if (!fp) {
printf("OPEN_FAIL %s\nSTATUS FAIL\n", outpath);
return 1;
}
hdr[0] = (uint32_t)nframes;
hdr[1] = WIDTH;
hdr[2] = HEIGHT;
hdr[3] = WIDTH * 4;
if (fwrite("R18FRM01", 1, 8, fp) != 8 ||
fwrite(hdr, 1, sizeof(hdr), fp) != sizeof(hdr) || fflush(fp)) {
printf("WRITE_FAIL header\nSTATUS FAIL\n");
fclose(fp);
return 1;
}
for (f = 0; f < nframes; f++) {
uchar4 *dst = d_fb;
uint64_t sum = 0;
size_t i;
t0 = now_ms();
plasma_kernel<<<blocks, 256>>>(dst, WIDTH, HEIGHT, f);
CHECK(cudaGetLastError());
CHECK(cudaDeviceSynchronize());
t1 = now_ms();
km += t1 - t0;
CHECK(cudaMemcpy(h_fb, dst, fbytes, cudaMemcpyDeviceToHost));
cm += now_ms() - t1;
for (i = 0; i < fbytes; i++)
sum += h_fb[i];
if (fwrite(h_fb, 1, fbytes, fp) != fbytes || fflush(fp)) {
printf("WRITE_FAIL frame %d\nSTATUS FAIL\n", f);
fclose(fp);
return 1;
}
printf("FRAME %03d fnv=0x%016llx mean=%u kms=%.1f cms=%.1f wall=%.1f\n",
f, (unsigned long long)fnv1a(h_fb, fbytes),
(unsigned)(sum / fbytes), km, cm, now_ms() - wall0);
}
fclose(fp);
printf("KERNEL frames=%d %dx%d blocks=%d launches=%d total_ms=%.2f per_frame_ms=%.2f\n",
nframes, WIDTH, HEIGHT, blocks, nframes, km, km / nframes);
printf("COPY frames=%d total_ms=%.2f per_frame_ms=%.2f\n", nframes,
cm, cm / nframes);
printf("WRITE path=%s bytes=%zu magic=R18FRM01 count=%d %dx%d stride=%d wall_ms=%.1f\n",
outpath, 24 + fbytes * (size_t)nframes, nframes, WIDTH, HEIGHT,
WIDTH * 4, now_ms() - wall0);
printf("STATUS OK\n");
CHECK(cudaFree(d_fb));
free(h_fb);
return 0;
}
|
