-
Notifications
You must be signed in to change notification settings - Fork 6
/
cltest.nt
219 lines (205 loc) · 8 KB
/
cltest.nt
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
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
218
219
module cltest;
import c.CL.cl;
import sys, std.string, std.util;
pragma(lib, "OpenCL");
extern(C) cl_context clCreateContextFromType(
cl_context_properties *properties,
cl_device_type device_type,
void function(char* errinfo, void* private_info, size_t cb, void* user_data) pfn_notify,
void* user_data, cl_int* errcode_ret);
extern(C) cl_int clBuildProgram(
cl_program, cl_uint, cl_device_id*, char*,
void function(cl_program, void*), void*);
template dgwrapper(T) {
extern(C) void callHolder(T t, void* ptr) {
auto trip = *(void*, void delegate(T) dg)*:ptr;
auto _threadlocal = getThreadlocal();
trip[1](t);
}
auto dgwrapper(void delegate(T) dg) {
auto ptr = new (void*, void delegate(T));
(*ptr) = (null, dg);
return (&callHolder, void*:ptr);
}
}
void clCheckRes (int i) {
if (i != 0) {
fail "CL failed with $i! ";
}
}
template clCheckCall(alias A) {
template clCheckCall(T) {
type-of A(value-of!T, null) clCheckCall(T t) {
int error;
onExit clCheckRes (error);
return A(t, &error);
}
}
}
cl_context createContextFromType(cl_context_properties[] props, cl_device_type type, void delegate(char* errinfo, void* private_info, size_t cb) notify) {
cl_int ret;
auto tup = dgwrapper!(char*, void*, size_t)(void delegate((char*,void*,size_t)):notify);
props ~= cl_context_properties:0;
return clCheckCall!clCreateContextFromType (props.ptr, type, (ParamTypes type-of &clCreateContextFromType)[2]: tup[0], tup[1]);
}
import sdl;
int main() {
auto openclSource = "
__kernel void mandel(__global int* res, int2 size, int iters, float4 rect) {
float blend = 0;
int aa = 4;
int ix = get_global_id(0), iy = get_global_id(1);
switch (aa) { case 4: ; }
for (int sy = 0; sy < aa; sy++) {
for (int sx = 0; sx < aa; sx++) {
float x = ix + (float) sx / aa, y = iy + (float) sy / aa;
float2 c = (float2) (x, y) / (float2) (size.x, size.y);
float2 rectsize = rect.zw - rect.xy;
c = rect.xy + c * rectsize;
float2 z = c;
int i;
for (i = 0; i < iters; ++i) {
float2 p = z * z;
if (p.x + p.y > 4) break;
z = (float2) (p.x - p.y, 2 * z.x * z.y) + c;
}
blend += i;
}
}
res[iy * size.x + ix] = (int) (blend / (aa * aa));
}"[];
int ids;
clCheckRes clGetPlatformIDs(0, null, &ids);
auto platforms = new cl_platform_id[] ids;
clCheckRes clGetPlatformIDs(ids, platforms.ptr, null);
writeln "$ids platform(s). ";
cl_device_id[] getDevices(cl_platform_id platf) {
int devs;
clCheckRes clGetDeviceIDs (platf, CL_DEVICE_TYPE_ALL, 0, null, &devs);
if (!devs) return null;
auto devlist = new cl_device_id[] devs;
clCheckRes clGetDeviceIDs (platf, CL_DEVICE_TYPE_ALL, devs, devlist.ptr, null);
return devlist;
}
while (int i, cl_platform_id platf) <- [for k <- 0..ids: (k, platforms[k])] {
auto things = [
("Profile"[], CL_PLATFORM_PROFILE),
("Version"[], CL_PLATFORM_VERSION),
("Name"[], CL_PLATFORM_NAME),
("Vendor"[], CL_PLATFORM_VENDOR),
("Extensions"[], CL_PLATFORM_EXTENSIONS)];
while (string info, int enum) <- things
{
int size;
clCheckRes clGetPlatformInfo (platf, enum, 0, null, &size);
auto store = new char[] size;
onExit store.free();
clCheckRes clGetPlatformInfo (platf, enum, size, store.ptr, int*:null);
writeln "$i: $info = $store";
}
auto devlist = getDevices(platf);
onExit devlist.free;
if (!devlist.length) {
writeln "No devices. ";
} else {
string devinfo; int enum2;
for (int k <- ints) && (auto dev <- devlist) {
alias things = [("Extensions"[], CL_DEVICE_EXTENSIONS),
("Name"[], CL_DEVICE_NAME),
("Profile"[], CL_DEVICE_PROFILE),
("Vendor"[], CL_DEVICE_VENDOR),
("Version"[], CL_DEVICE_VERSION),
("DriverVersion"[], CL_DRIVER_VERSION)];
for (string devinfo, int enum2) <- things {
int size;
clCheckRes clGetDeviceInfo (dev, enum2, 0, null, &size);
scope devstore = new char[] size;
clCheckRes clGetDeviceInfo (dev, enum2, size, devstore.ptr, int*:null);
writeln " $k: $devinfo = $devstore ($size)";
}
}
}
}
auto platf = platforms[0];
int device = 0;
cl_context_properties[] props;
props ~= CL_CONTEXT_PLATFORM;
props ~= cl_context_properties: platf;
auto ctx = createContextFromType(props, CL_DEVICE_TYPE_ALL, null);
onExit clReleaseContext (ctx);
writeln "Context created. ";
auto dev = getDevices(platf)[device];
auto queue = clCheckCall!clCreateCommandQueue (ctx, dev, 0);
writeln "Command queue created. ";
auto rect = vec4f(-2, -2, 2, 2);
auto iters = cl_int:512, size = (cl_int:800, cl_int:600), output = new int[] (size[0]*size[1]);
auto outvec = clCheckCall!clCreateBuffer (ctx, CL_MEM_WRITE_ONLY,
(size-of int) * size[0] * size[1], null);
onExit clReleaseMemObject (outvec);
writeln "Buffers created. ";
auto sourcelines = [for line <- splitAt(once openclSource, "\n"): line ~ "\n\x00"].eval[];
writeln "$(sourcelines.length) lines of source. ";
auto prog = clCreateProgramWithSource(ctx, sourcelines.length,
[for line <- sourcelines: line.ptr].eval.ptr, null, null);
onExit clReleaseProgram (prog);
writeln "Building. ";
{
if auto err = clBuildProgram (prog, 0, null x 4) {
int len;
clGetProgramBuildInfo (prog, dev, CL_PROGRAM_BUILD_LOG, 0, null, &len);
auto str = new char[] len;
clGetProgramBuildInfo (prog, dev, CL_PROGRAM_BUILD_LOG, len, str.ptr, null);
writeln "Failed to build: $str";
_interrupt 3;
}
}
writeln "Program built. ";
auto addKernel = clCheckCall!clCreateKernel (prog, "mandel".ptr);
onExit clReleaseKernel (addKernel);
writeln "Kernel created. ";
void calc() {
clCheckRes clSetKernelArg (addKernel, 0, size-of type-of outvec, void*:&outvec);
clCheckRes clSetKernelArg (addKernel, 1, size-of type-of size, void*:&size);
clCheckRes clSetKernelArg (addKernel, 2, size-of int, void*:&iters);
clCheckRes clSetKernelArg (addKernel, 3, size-of type-of rect, void*:&rect);
auto workSize = [size[0], size[1]];
clCheckRes clEnqueueNDRangeKernel (queue, addKernel, 2, null, workSize.ptr, null, 0, null, null);
// read-back
clCheckRes clEnqueueReadBuffer (queue, outvec, CL_TRUE, 0, (size-of int) * size[0] * size[1], output.ptr, 0, null, null);
}
SDL_Init(32); // video
auto surface = SDL_Surface*: SDL_SetVideoMode(size, 0, SDL_ANYFORMAT);
bool update() {
SDL_Flip(surface);
SDL_Event ev;
while SDL_PollEvent(&ev) {
if ev.type == SDL_KEYDOWN {
auto sym = (*(SDL_KeyboardEvent*:&ev)).keysym.sym;
using rect {
if (sym == 273) { auto dist = (w - y) * 0.1; y -= dist; w -= dist; }
else if (sym == 274) { auto dist = (w - y) * 0.1; y += dist; w += dist; }
else if (sym == 275) { auto dist = (z - x) * 0.1; x += dist; z += dist; }
else if (sym == 276) { auto dist = (z - x) * 0.1; x -= dist; z -= dist; }
else if (sym == 43) { auto dist = zw - xy, half = 0.5 * (xy + zw); dist *= 0.8; (x,y) = half - dist * 0.5; (z,w) = half + dist * 0.5; }
else if (sym == 45) { auto dist = zw - xy, half = 0.5 * (xy + zw); dist /= 0.8; (x,y) = half - dist * 0.5; (z,w) = half + dist * 0.5; }
else writeln "Key $(int:sym)";
}
}
if ev.type == SDL_QUIT { exit(0); return true; } // QUIT
}
return false;
}
void draw() {
int factor1 = 255, factor2 = 256 * 255, factor3 = 256 * 256 * 255;
auto f1f = float:factor1, f2f = float:factor2, f3f = float:factor3;
for (int y = 0; y < size[1]; y++) {
auto p = &((int*:surface.pixels)[y * int:surface.w]);
for (int x = 0; x < size[0]; x++) {
auto f = vec3f(output[y * size[0] + x] / 512.0);
*(p++) = int:(f1f * f[2]) + int:(f2f * f[1]) & factor2 + int:(f3f * f[0]) & factor3;
}
}
}
while !update() { calc(); draw(); }
return 0;
}