-
Notifications
You must be signed in to change notification settings - Fork 0
/
Copy pathcosmos_render.cpp
363 lines (216 loc) · 7.86 KB
/
cosmos_render.cpp
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
220
221
222
223
224
225
226
227
228
229
230
231
232
233
234
235
236
237
238
239
240
241
242
243
244
245
246
247
248
249
250
251
252
253
254
255
256
257
258
259
260
261
262
263
264
265
266
267
268
269
270
271
272
273
274
275
276
277
278
279
280
281
282
283
284
285
286
287
288
289
290
291
292
293
294
295
296
297
298
299
300
301
302
303
304
305
306
307
308
309
310
311
312
313
314
315
316
317
318
319
320
321
322
323
324
325
326
327
328
329
330
331
332
333
334
335
336
337
338
339
340
341
342
343
344
345
346
347
348
349
350
351
352
353
354
355
356
357
358
359
360
361
362
363
/*
### cosmos_render
To compile, use the command
```bash
clang++ cosmos_render.cpp -o cosmos_render -std=c++11 -lOpenCL
```
To run, use the command
```bash
./cosmos_render <frames> [bodies=24576]
```
This will cause cosmos_render to output a bunch of rendered frames to the
directory that it was started in. These rendered frames will be image files,
produced by rendering the corresponding simulation frames generated by running
cosmos_simulate.
## Even more speed
To speed up simulations and renders even more, compile with
```bash
clang++ cosmos_tool.cpp -o cosmos_tool -std=c++11 -lOpenCL -Ofast -march=native
```
*/
#include <iostream>
#include <sstream>
#include <fstream>
#include <ctime>
#include <memory>
// Include stb_image_write.
#define STB_IMAGE_WRITE_IMPLEMENTATION
#include "stb_image_write.h"
// Include OpenCL.
#ifdef __APPLE__
#define CL_SILENCE_DEPRECATION
#include <OpenCL/OpenCL.h>
#else
#include <CL/cl.h>
#endif
// Include the kernel source.
#define __stringify(source) #source
const char* kernel_source =
#include "cosmos_render.cl"
#undef __stringify
size_t kernel_source_size = strlen(kernel_source);
// Include the thermal colormap.
#define THERMAL_OPEN_CL
#include "thermal_colormap.h"
// Write a message to std::cout.
void say(std::string message)
{
std::cout << message << std::endl;
}
// Entry point.
int main(int argc, char** argv)
{
// Parse command line arguments.
if (argc != 2)
{
std::cout << "Usage: " << argv[0] << " <frames> [bodies=24576]" << std::endl;
return EXIT_FAILURE;
}
int frames = atoi(argv[1]);
int n = 24576;
if (argc == 3)
{
n = atoi(argv[2]);
}
// Create variables to hold return codes.
cl_int r_code;
cl_int r_code1;
cl_int r_code2;
// Create identifier objects to hold information about the available
// platforms and available devices.
cl_platform_id platform_id = NULL;
cl_device_id device_id = NULL;
// Create unsigned integer objects to hold the amount of available
// platforms and available devices.
cl_uint num_platforms;
cl_uint num_devices;
// Get the first available platform and store the amount of available
// platforms.
clGetPlatformIDs(1, &platform_id, &num_platforms);
// Get the first available device on the first available platform. Store
// the amount of available devices. This device will be referred to as the
// 'default device'.
clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_DEFAULT, 1, &device_id, &num_devices);
// Create an OpenCL context on the default device.
cl_context context = clCreateContext(0, 1, &device_id, NULL, NULL, &r_code);
// Make sure the OpenCL context was created successfully.
if (r_code != CL_SUCCESS)
{
say("Could not create an OpenCL context.");
return EXIT_FAILURE;
}
// Create an OpenCL command queue.
cl_command_queue command_queue = clCreateCommandQueue(context, device_id, 0, &r_code);
// Make sure the OpenCL command queue was created successfully.
if (r_code != CL_SUCCESS)
{
say("Could not create an OpenCL command queue.");
return EXIT_FAILURE;
}
// Allocate a buffer to hold the frame body data, CPU side.
cl_float4* cpu_bodies = (cl_float4*)malloc(n * sizeof(cl_float4));
// Allocate a buffer to hold the frame body data, GPU side.
cl_mem gpu_bodies = clCreateBuffer(context, CL_MEM_READ_WRITE, n * sizeof(cl_float4), NULL, &r_code);
if (r_code != CL_SUCCESS)
{
say("Could not allocate GPU memory.");
return EXIT_FAILURE;
}
// PARAM: The dimensions of the image (which is square).
int res = 800;
int stride_in_bytes = res * 4;
// Allocate a buffer to hold the rendered frame, CPU side.
unsigned char* cpu_img = (unsigned char*)malloc(res * res * 4 * sizeof(unsigned char));
// Allocate a buffer to hold the rendered frame, GPU side.
cl_mem gpu_img = clCreateBuffer(context, CL_MEM_READ_WRITE, res * res * sizeof(cl_uint), NULL, &r_code);
if (r_code != CL_SUCCESS)
{
say("Could not allocate GPU memory.");
return EXIT_FAILURE;
}
// Allocate an array to hold the thermal colormap, GPU side.
cl_mem gpu_colormap = clCreateBuffer(context, CL_MEM_READ_WRITE, 256 * sizeof(cl_uint), NULL, &r_code);
if (r_code != CL_SUCCESS)
{
say("Could not allocate GPU memory.");
return EXIT_FAILURE;
}
// Copy the thermal colormap from the CPU to the GPU.
r_code = clEnqueueWriteBuffer(command_queue, gpu_colormap, CL_TRUE, 0, 256 * sizeof(cl_uint), thermal_colormap, 0, NULL, NULL);
if (r_code != CL_SUCCESS)
{
say("Could not copy the thermal colormap array from the CPU to the GPU.");
return EXIT_FAILURE;
}
// Create an OpenCL program from the kernel source.
cl_program program = clCreateProgramWithSource(context, 1, (const char**)&kernel_source, (const size_t*)&kernel_source_size, &r_code);
// Make sure the OpenCL program was created successfully.
if (r_code != CL_SUCCESS)
{
say("Could not create an OpenCL program.");
return EXIT_FAILURE;
}
// Build the OpenCL program.
r_code = clBuildProgram(program, 1, &device_id, NULL, NULL, NULL);
// Make sure the OpenCL program was built successfully.
if (r_code != CL_SUCCESS)
{
say("Could not build an OpenCL program.");
return EXIT_FAILURE;
}
// Create the OpenCL kernel from the function "n_body_render" within the
// OpenCL program.
cl_kernel kernel = clCreateKernel(program, "n_body_render", &r_code);
// Make sure the OpenCL kernel was created successfully.
if (r_code != CL_SUCCESS)
{
say("Could not create an OpenCL kernel.");
return EXIT_FAILURE;
}
// Set the n parameter of the OpenCL kernel.
clSetKernelArg(kernel, 1, sizeof(cl_int), &n);
// Set the res parameter of the OpenCL kernel.
clSetKernelArg(kernel, 2, sizeof(cl_int), &res);
// Set the state parameter of the OpenCL kernel.
clSetKernelArg(kernel, 3, sizeof(cl_mem), (void*)&gpu_bodies);
// Set the output parameter of the OpenCL kernel.
clSetKernelArg(kernel, 4, sizeof(cl_mem), (void*)&gpu_img);
// Set the thermal_colormap parameter of the OpenCL kernel.
clSetKernelArg(kernel, 5, sizeof(cl_mem), (void*)&gpu_colormap);
// PARAM: local_work_size should be modified depending on your GPU. This
// should be tested on a realtime renderer first, such as
// CobaltXII/boiler/experimental/n_body_cl/.
size_t global_work_size = res * res;
size_t local_work_size = 256;
// Start the render!
for (int i = 0; i < frames; i++)
{
// Get the starting time.
clock_t frame_start = clock();
// Load the current frame.
std::stringstream name_builder;
name_builder << "frame_" << i << ".dat";
std::ifstream frame(name_builder.str());
frame.read((char*)cpu_bodies, n * sizeof(cl_float4));
frame.close();
// Copy the bodies array from the CPU to the GPU.
r_code = clEnqueueWriteBuffer(command_queue, gpu_bodies, CL_TRUE, 0, n * sizeof(cl_float4), cpu_bodies, 0, NULL, NULL);
if (r_code != CL_SUCCESS)
{
say("Could not copy the body array from the CPU to the GPU.");
return EXIT_FAILURE;
}
// PARAM: The zoom factor.
float inv_scale = i / 60.0f;
// Set the inv_scale parameter of the OpenCL kernel.
clSetKernelArg(kernel, 0, sizeof(cl_float), &inv_scale);
// Generate the current frame.
r_code = clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL, &global_work_size, &local_work_size, 0, NULL, NULL);
if (r_code != CL_SUCCESS)
{
say("Could not invoke the render kernel.");
return EXIT_FAILURE;
}
// Read output back into local CPU memory.
clEnqueueReadBuffer(command_queue, gpu_img, CL_TRUE, 0, res * res * sizeof(cl_uint), cpu_img, 0, NULL, NULL);
// Save the image.
std::stringstream out_name_builder;
out_name_builder << "render_" << i << ".png";
stbi_write_png(out_name_builder.str().c_str(), res, res, 4, cpu_img, stride_in_bytes);
// Get the ending time.
clock_t frame_end = clock();
// Print the frame elapsed time.
std::cout << "Frame " << i << " rendered in " << float(frame_end - frame_start) / float(CLOCKS_PER_SEC) << " s" << std::endl;
}
return EXIT_SUCCESS;
}