Cycles OpenCL: a few fixes to get things compiling after kernel changes,
[blender.git] / intern / cycles / device / device_opencl.cpp
1 /*
2  * Copyright 2011, Blender Foundation.
3  *
4  * This program is free software; you can redistribute it and/or
5  * modify it under the terms of the GNU General Public License
6  * as published by the Free Software Foundation; either version 2
7  * of the License, or (at your option) any later version.
8  *
9  * This program is distributed in the hope that it will be useful,
10  * but WITHOUT ANY WARRANTY; without even the implied warranty of
11  * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
12  * GNU General Public License for more details.
13  *
14  * You should have received a copy of the GNU General Public License
15  * along with this program; if not, write to the Free Software Foundation,
16  * Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
17  */
18
19 #ifdef WITH_OPENCL
20
21 #include <stdio.h>
22 #include <stdlib.h>
23 #include <string.h>
24
25 #include "device.h"
26 #include "device_intern.h"
27
28 #include "buffers.h"
29
30 #include "util_foreach.h"
31 #include "util_map.h"
32 #include "util_math.h"
33 #include "util_md5.h"
34 #include "util_opencl.h"
35 #include "util_opengl.h"
36 #include "util_path.h"
37 #include "util_time.h"
38
39 CCL_NAMESPACE_BEGIN
40
41 #define CL_MEM_PTR(p) ((cl_mem)(unsigned long)(p))
42
43 static cl_device_type opencl_device_type()
44 {
45         char *device = getenv("CYCLES_OPENCL_TEST");
46
47         if(device) {
48                 if(strcmp(device, "ALL") == 0)
49                         return CL_DEVICE_TYPE_ALL;
50                 else if(strcmp(device, "DEFAULT") == 0)
51                         return CL_DEVICE_TYPE_DEFAULT;
52                 else if(strcmp(device, "CPU") == 0)
53                         return CL_DEVICE_TYPE_CPU;
54                 else if(strcmp(device, "GPU") == 0)
55                         return CL_DEVICE_TYPE_GPU;
56                 else if(strcmp(device, "ACCELERATOR") == 0)
57                         return CL_DEVICE_TYPE_ACCELERATOR;
58         }
59
60         return CL_DEVICE_TYPE_GPU|CL_DEVICE_TYPE_ACCELERATOR;
61 }
62
63 class OpenCLDevice : public Device
64 {
65 public:
66         TaskPool task_pool;
67         cl_context cxContext;
68         cl_command_queue cqCommandQueue;
69         cl_platform_id cpPlatform;
70         cl_device_id cdDevice;
71         cl_program cpProgram;
72         cl_kernel ckPathTraceKernel;
73         cl_kernel ckFilmConvertKernel;
74         cl_int ciErr;
75         map<string, device_vector<uchar>*> const_mem_map;
76         map<string, device_memory*> mem_map;
77         device_ptr null_mem;
78         bool device_initialized;
79         string platform_name;
80
81         const char *opencl_error_string(cl_int err)
82         {
83                 switch (err) {
84                         case CL_SUCCESS: return "Success!";
85                         case CL_DEVICE_NOT_FOUND: return "Device not found.";
86                         case CL_DEVICE_NOT_AVAILABLE: return "Device not available";
87                         case CL_COMPILER_NOT_AVAILABLE: return "Compiler not available";
88                         case CL_MEM_OBJECT_ALLOCATION_FAILURE: return "Memory object allocation failure";
89                         case CL_OUT_OF_RESOURCES: return "Out of resources";
90                         case CL_OUT_OF_HOST_MEMORY: return "Out of host memory";
91                         case CL_PROFILING_INFO_NOT_AVAILABLE: return "Profiling information not available";
92                         case CL_MEM_COPY_OVERLAP: return "Memory copy overlap";
93                         case CL_IMAGE_FORMAT_MISMATCH: return "Image format mismatch";
94                         case CL_IMAGE_FORMAT_NOT_SUPPORTED: return "Image format not supported";
95                         case CL_BUILD_PROGRAM_FAILURE: return "Program build failure";
96                         case CL_MAP_FAILURE: return "Map failure";
97                         case CL_INVALID_VALUE: return "Invalid value";
98                         case CL_INVALID_DEVICE_TYPE: return "Invalid device type";
99                         case CL_INVALID_PLATFORM: return "Invalid platform";
100                         case CL_INVALID_DEVICE: return "Invalid device";
101                         case CL_INVALID_CONTEXT: return "Invalid context";
102                         case CL_INVALID_QUEUE_PROPERTIES: return "Invalid queue properties";
103                         case CL_INVALID_COMMAND_QUEUE: return "Invalid command queue";
104                         case CL_INVALID_HOST_PTR: return "Invalid host pointer";
105                         case CL_INVALID_MEM_OBJECT: return "Invalid memory object";
106                         case CL_INVALID_IMAGE_FORMAT_DESCRIPTOR: return "Invalid image format descriptor";
107                         case CL_INVALID_IMAGE_SIZE: return "Invalid image size";
108                         case CL_INVALID_SAMPLER: return "Invalid sampler";
109                         case CL_INVALID_BINARY: return "Invalid binary";
110                         case CL_INVALID_BUILD_OPTIONS: return "Invalid build options";
111                         case CL_INVALID_PROGRAM: return "Invalid program";
112                         case CL_INVALID_PROGRAM_EXECUTABLE: return "Invalid program executable";
113                         case CL_INVALID_KERNEL_NAME: return "Invalid kernel name";
114                         case CL_INVALID_KERNEL_DEFINITION: return "Invalid kernel definition";
115                         case CL_INVALID_KERNEL: return "Invalid kernel";
116                         case CL_INVALID_ARG_INDEX: return "Invalid argument index";
117                         case CL_INVALID_ARG_VALUE: return "Invalid argument value";
118                         case CL_INVALID_ARG_SIZE: return "Invalid argument size";
119                         case CL_INVALID_KERNEL_ARGS: return "Invalid kernel arguments";
120                         case CL_INVALID_WORK_DIMENSION: return "Invalid work dimension";
121                         case CL_INVALID_WORK_GROUP_SIZE: return "Invalid work group size";
122                         case CL_INVALID_WORK_ITEM_SIZE: return "Invalid work item size";
123                         case CL_INVALID_GLOBAL_OFFSET: return "Invalid global offset";
124                         case CL_INVALID_EVENT_WAIT_LIST: return "Invalid event wait list";
125                         case CL_INVALID_EVENT: return "Invalid event";
126                         case CL_INVALID_OPERATION: return "Invalid operation";
127                         case CL_INVALID_GL_OBJECT: return "Invalid OpenGL object";
128                         case CL_INVALID_BUFFER_SIZE: return "Invalid buffer size";
129                         case CL_INVALID_MIP_LEVEL: return "Invalid mip-map level";
130                         default: return "Unknown";
131                 }
132         }
133
134         bool opencl_error(cl_int err)
135         {
136                 if(err != CL_SUCCESS) {
137                         string message = string_printf("OpenCL error (%d): %s", err, opencl_error_string(err));
138                         if(error_msg == "")
139                                 error_msg = message;
140                         fprintf(stderr, "%s\n", message.c_str());
141                         return true;
142                 }
143
144                 return false;
145         }
146
147         void opencl_error(const string& message)
148         {
149                 if(error_msg == "")
150                         error_msg = message;
151                 fprintf(stderr, "%s\n", message.c_str());
152         }
153
154         void opencl_assert(cl_int err)
155         {
156                 if(err != CL_SUCCESS) {
157                         string message = string_printf("OpenCL error (%d): %s", err, opencl_error_string(err));
158                         if(error_msg == "")
159                                 error_msg = message;
160                         fprintf(stderr, "%s\n", message.c_str());
161 #ifndef NDEBUG
162                         abort();
163 #endif
164                 }
165         }
166
167         OpenCLDevice(DeviceInfo& info, Stats &stats, bool background_)
168           : Device(stats)
169         {
170                 background = background_;
171                 cpPlatform = NULL;
172                 cxContext = NULL;
173                 cqCommandQueue = NULL;
174                 cpProgram = NULL;
175                 ckPathTraceKernel = NULL;
176                 ckFilmConvertKernel = NULL;
177                 null_mem = 0;
178                 device_initialized = false;
179
180                 /* setup platform */
181                 cl_uint num_platforms;
182
183                 ciErr = clGetPlatformIDs(0, NULL, &num_platforms);
184                 if(opencl_error(ciErr))
185                         return;
186
187                 if(num_platforms == 0) {
188                         opencl_error("OpenCL: no platforms found.");
189                         return;
190                 }
191
192                 ciErr = clGetPlatformIDs(1, &cpPlatform, NULL);
193                 if(opencl_error(ciErr))
194                         return;
195
196                 char name[256];
197                 clGetPlatformInfo(cpPlatform, CL_PLATFORM_NAME, sizeof(name), &name, NULL);
198                 platform_name = name;
199
200                 /* get devices */
201                 vector<cl_device_id> device_ids;
202                 cl_uint num_devices;
203
204                 if(opencl_error(clGetDeviceIDs(cpPlatform, opencl_device_type(), 0, NULL, &num_devices)))
205                         return;
206
207                 if(info.num > num_devices) {
208                         if(num_devices == 0)
209                                 opencl_error("OpenCL: no devices found.");
210                         else
211                                 opencl_error("OpenCL: specified device not found.");
212                         return;
213                 }
214
215                 device_ids.resize(num_devices);
216                 
217                 if(opencl_error(clGetDeviceIDs(cpPlatform, opencl_device_type(), num_devices, &device_ids[0], NULL)))
218                         return;
219
220                 cdDevice = device_ids[info.num];
221
222                 /* create context */
223                 cxContext = clCreateContext(0, 1, &cdDevice, NULL, NULL, &ciErr);
224                 if(opencl_error(ciErr))
225                         return;
226
227                 cqCommandQueue = clCreateCommandQueue(cxContext, cdDevice, 0, &ciErr);
228                 if(opencl_error(ciErr))
229                         return;
230
231                 null_mem = (device_ptr)clCreateBuffer(cxContext, CL_MEM_READ_ONLY, 1, NULL, &ciErr);
232                 device_initialized = true;
233         }
234
235         bool opencl_version_check()
236         {
237                 char version[256];
238
239                 int major, minor, req_major = 1, req_minor = 0;
240
241                 clGetPlatformInfo(cpPlatform, CL_PLATFORM_VERSION, sizeof(version), &version, NULL);
242
243                 if(sscanf(version, "OpenCL %d.%d", &major, &minor) < 2) {
244                         opencl_error(string_printf("OpenCL: failed to parse platform version string (%s).", version));
245                         return false;
246                 }
247
248                 if(!((major == req_major && minor >= req_minor) || (major > req_major))) {
249                         opencl_error(string_printf("OpenCL: platform version 1.1 or later required, found %d.%d", major, minor));
250                         return false;
251                 }
252
253                 clGetDeviceInfo(cdDevice, CL_DEVICE_OPENCL_C_VERSION, sizeof(version), &version, NULL);
254
255                 if(sscanf(version, "OpenCL C %d.%d", &major, &minor) < 2) {
256                         opencl_error(string_printf("OpenCL: failed to parse OpenCL C version string (%s).", version));
257                         return false;
258                 }
259
260                 if(!((major == req_major && minor >= req_minor) || (major > req_major))) {
261                         opencl_error(string_printf("OpenCL: C version 1.1 or later required, found %d.%d", major, minor));
262                         return false;
263                 }
264
265                 /* we don't check CL_DEVICE_VERSION since for e.g. nvidia sm 1.3 cards this is
266                  * 1.0 even if the language features are there, just limited shared memory */
267
268                 return true;
269         }
270
271         bool load_binary(const string& kernel_path, const string& clbin)
272         {
273                 /* read binary into memory */
274                 vector<uint8_t> binary;
275
276                 if(!path_read_binary(clbin, binary)) {
277                         opencl_error(string_printf("OpenCL failed to read cached binary %s.", clbin.c_str()));
278                         return false;
279                 }
280
281                 /* create program */
282                 cl_int status;
283                 size_t size = binary.size();
284                 const uint8_t *bytes = &binary[0];
285
286                 cpProgram = clCreateProgramWithBinary(cxContext, 1, &cdDevice,
287                         &size, &bytes, &status, &ciErr);
288
289                 if(opencl_error(status) || opencl_error(ciErr)) {
290                         opencl_error(string_printf("OpenCL failed create program from cached binary %s.", clbin.c_str()));
291                         return false;
292                 }
293
294                 if(!build_kernel(kernel_path))
295                         return false;
296
297                 return true;
298         }
299
300         bool save_binary(const string& clbin)
301         {
302                 size_t size = 0;
303                 clGetProgramInfo(cpProgram, CL_PROGRAM_BINARY_SIZES, sizeof(size_t), &size, NULL);
304
305                 if(!size)
306                         return false;
307
308                 vector<uint8_t> binary(size);
309                 uint8_t *bytes = &binary[0];
310
311                 clGetProgramInfo(cpProgram, CL_PROGRAM_BINARIES, sizeof(uint8_t*), &bytes, NULL);
312
313                 if(!path_write_binary(clbin, binary)) {
314                         opencl_error(string_printf("OpenCL failed to write cached binary %s.", clbin.c_str()));
315                         return false;
316                 }
317
318                 return true;
319         }
320
321         string kernel_build_options()
322         {
323                 string build_options = " -cl-fast-relaxed-math ";
324                 
325                 if(platform_name == "NVIDIA CUDA")
326                         build_options += "-D__KERNEL_OPENCL_NVIDIA__ -cl-nv-maxrregcount=24 -cl-nv-verbose ";
327
328                 else if(platform_name == "Apple")
329                         build_options += "-D__KERNEL_OPENCL_APPLE__ -Wno-missing-prototypes";
330
331                 else if(platform_name == "AMD Accelerated Parallel Processing")
332                         build_options += "-D__KERNEL_OPENCL_AMD__ ";
333
334                 return build_options;
335         }
336
337         bool build_kernel(const string& kernel_path)
338         {
339                 string build_options = kernel_build_options();
340         
341                 ciErr = clBuildProgram(cpProgram, 0, NULL, build_options.c_str(), NULL, NULL);
342
343                 if(ciErr != CL_SUCCESS) {
344                         /* show build errors */
345                         char *build_log;
346                         size_t ret_val_size;
347
348                         clGetProgramBuildInfo(cpProgram, cdDevice, CL_PROGRAM_BUILD_LOG, 0, NULL, &ret_val_size);
349
350                         build_log = new char[ret_val_size+1];
351                         clGetProgramBuildInfo(cpProgram, cdDevice, CL_PROGRAM_BUILD_LOG, ret_val_size, build_log, NULL);
352
353                         build_log[ret_val_size] = '\0';
354                         opencl_error("OpenCL build failed: errors in console");
355                         fprintf(stderr, "%s\n", build_log);
356
357                         delete[] build_log;
358
359                         return false;
360                 }
361
362                 return true;
363         }
364
365         bool compile_kernel(const string& kernel_path, const string& kernel_md5)
366         {
367                 /* we compile kernels consisting of many files. unfortunately opencl
368                  * kernel caches do not seem to recognize changes in included files.
369                  * so we force recompile on changes by adding the md5 hash of all files */
370                 string source = "#include \"kernel.cl\" // " + kernel_md5 + "\n";
371                 source = path_source_replace_includes(source, kernel_path);
372
373                 size_t source_len = source.size();
374                 const char *source_str = source.c_str();
375
376                 cpProgram = clCreateProgramWithSource(cxContext, 1, &source_str, &source_len, &ciErr);
377
378                 if(opencl_error(ciErr))
379                         return false;
380
381                 double starttime = time_dt();
382                 printf("Compiling OpenCL kernel ...\n");
383
384                 if(!build_kernel(kernel_path))
385                         return false;
386
387                 printf("Kernel compilation finished in %.2lfs.\n", time_dt() - starttime);
388
389                 return true;
390         }
391
392         string device_md5_hash()
393         {
394                 MD5Hash md5;
395                 char version[256], driver[256], name[256], vendor[256];
396
397                 clGetPlatformInfo(cpPlatform, CL_PLATFORM_VENDOR, sizeof(vendor), &vendor, NULL);
398                 clGetDeviceInfo(cdDevice, CL_DEVICE_VERSION, sizeof(version), &version, NULL);
399                 clGetDeviceInfo(cdDevice, CL_DEVICE_NAME, sizeof(name), &name, NULL);
400                 clGetDeviceInfo(cdDevice, CL_DRIVER_VERSION, sizeof(driver), &driver, NULL);
401
402                 md5.append((uint8_t*)vendor, strlen(vendor));
403                 md5.append((uint8_t*)version, strlen(version));
404                 md5.append((uint8_t*)name, strlen(name));
405                 md5.append((uint8_t*)driver, strlen(driver));
406
407                 string options = kernel_build_options();
408                 md5.append((uint8_t*)options.c_str(), options.size());
409
410                 return md5.get_hex();
411         }
412
413         bool load_kernels(bool experimental)
414         {
415                 /* verify if device was initialized */
416                 if(!device_initialized) {
417                         fprintf(stderr, "OpenCL: failed to initialize device.\n");
418                         return false;
419                 }
420
421                 /* verify we have right opencl version */
422                 if(!opencl_version_check())
423                         return false;
424
425                 /* md5 hash to detect changes */
426                 string kernel_path = path_get("kernel");
427                 string kernel_md5 = path_files_md5_hash(kernel_path);
428                 string device_md5 = device_md5_hash();
429
430                 /* try to use cache binary */
431                 string clbin = string_printf("cycles_kernel_%s_%s.clbin", device_md5.c_str(), kernel_md5.c_str());
432                 clbin = path_user_get(path_join("cache", clbin));
433
434                 if(path_exists(clbin)) {
435                         /* if exists already, try use it */
436                         if(!load_binary(kernel_path, clbin))
437                                 return false;
438                 }
439                 else {
440                         /* compile kernel */
441                         if(!compile_kernel(kernel_path, kernel_md5))
442                                 return false;
443
444                         /* save binary for reuse */
445                         save_binary(clbin);
446                 }
447
448                 /* find kernels */
449                 ckPathTraceKernel = clCreateKernel(cpProgram, "kernel_ocl_path_trace", &ciErr);
450                 if(opencl_error(ciErr))
451                         return false;
452
453                 ckFilmConvertKernel = clCreateKernel(cpProgram, "kernel_ocl_tonemap", &ciErr);
454                 if(opencl_error(ciErr))
455                         return false;
456
457                 return true;
458         }
459
460         ~OpenCLDevice()
461         {
462                 task_pool.stop();
463
464                 if(null_mem)
465                         clReleaseMemObject(CL_MEM_PTR(null_mem));
466
467                 map<string, device_vector<uchar>*>::iterator mt;
468                 for(mt = const_mem_map.begin(); mt != const_mem_map.end(); mt++) {
469                         mem_free(*(mt->second));
470                         delete mt->second;
471                 }
472
473                 if(ckPathTraceKernel)
474                         clReleaseKernel(ckPathTraceKernel);  
475                 if(ckFilmConvertKernel)
476                         clReleaseKernel(ckFilmConvertKernel);  
477                 if(cpProgram)
478                         clReleaseProgram(cpProgram);
479                 if(cqCommandQueue)
480                         clReleaseCommandQueue(cqCommandQueue);
481                 if(cxContext)
482                         clReleaseContext(cxContext);
483         }
484
485         void mem_alloc(device_memory& mem, MemoryType type)
486         {
487                 size_t size = mem.memory_size();
488
489                 if(type == MEM_READ_ONLY)
490                         mem.device_pointer = (device_ptr)clCreateBuffer(cxContext, CL_MEM_READ_ONLY, size, NULL, &ciErr);
491                 else if(type == MEM_WRITE_ONLY)
492                         mem.device_pointer = (device_ptr)clCreateBuffer(cxContext, CL_MEM_WRITE_ONLY, size, NULL, &ciErr);
493                 else
494                         mem.device_pointer = (device_ptr)clCreateBuffer(cxContext, CL_MEM_READ_WRITE, size, NULL, &ciErr);
495
496                 opencl_assert(ciErr);
497
498                 stats.mem_alloc(size);
499         }
500
501         void mem_copy_to(device_memory& mem)
502         {
503                 /* this is blocking */
504                 size_t size = mem.memory_size();
505                 ciErr = clEnqueueWriteBuffer(cqCommandQueue, CL_MEM_PTR(mem.device_pointer), CL_TRUE, 0, size, (void*)mem.data_pointer, 0, NULL, NULL);
506                 opencl_assert(ciErr);
507         }
508
509         void mem_copy_from(device_memory& mem, int y, int w, int h, int elem)
510         {
511                 size_t offset = elem*y*w;
512                 size_t size = elem*w*h;
513
514                 ciErr = clEnqueueReadBuffer(cqCommandQueue, CL_MEM_PTR(mem.device_pointer), CL_TRUE, offset, size, (uchar*)mem.data_pointer + offset, 0, NULL, NULL);
515                 opencl_assert(ciErr);
516         }
517
518         void mem_zero(device_memory& mem)
519         {
520                 if(mem.device_pointer) {
521                         memset((void*)mem.data_pointer, 0, mem.memory_size());
522                         mem_copy_to(mem);
523                 }
524         }
525
526         void mem_free(device_memory& mem)
527         {
528                 if(mem.device_pointer) {
529                         ciErr = clReleaseMemObject(CL_MEM_PTR(mem.device_pointer));
530                         mem.device_pointer = 0;
531                         opencl_assert(ciErr);
532
533                         stats.mem_free(mem.memory_size());
534                 }
535         }
536
537         void const_copy_to(const char *name, void *host, size_t size)
538         {
539                 if(const_mem_map.find(name) == const_mem_map.end()) {
540                         device_vector<uchar> *data = new device_vector<uchar>();
541                         data->copy((uchar*)host, size);
542
543                         mem_alloc(*data, MEM_READ_ONLY);
544                         const_mem_map[name] = data;
545                 }
546                 else {
547                         device_vector<uchar> *data = const_mem_map[name];
548                         data->copy((uchar*)host, size);
549                 }
550
551                 mem_copy_to(*const_mem_map[name]);
552         }
553
554         void tex_alloc(const char *name, device_memory& mem, bool interpolation, bool periodic)
555         {
556                 mem_alloc(mem, MEM_READ_ONLY);
557                 mem_copy_to(mem);
558                 mem_map[name] = &mem;
559         }
560
561         void tex_free(device_memory& mem)
562         {
563                 if(mem.data_pointer)
564                         mem_free(mem);
565         }
566
567         size_t global_size_round_up(int group_size, int global_size)
568         {
569                 int r = global_size % group_size;
570                 return global_size + ((r == 0)? 0: group_size - r);
571         }
572
573         void path_trace(RenderTile& rtile, int sample)
574         {
575                 /* cast arguments to cl types */
576                 cl_mem d_data = CL_MEM_PTR(const_mem_map["__data"]->device_pointer);
577                 cl_mem d_buffer = CL_MEM_PTR(rtile.buffer);
578                 cl_mem d_rng_state = CL_MEM_PTR(rtile.rng_state);
579                 cl_int d_x = rtile.x;
580                 cl_int d_y = rtile.y;
581                 cl_int d_w = rtile.w;
582                 cl_int d_h = rtile.h;
583                 cl_int d_sample = sample;
584                 cl_int d_offset = rtile.offset;
585                 cl_int d_stride = rtile.stride;
586
587                 /* sample arguments */
588                 int narg = 0;
589                 ciErr = 0;
590
591                 ciErr |= clSetKernelArg(ckPathTraceKernel, narg++, sizeof(d_data), (void*)&d_data);
592                 ciErr |= clSetKernelArg(ckPathTraceKernel, narg++, sizeof(d_buffer), (void*)&d_buffer);
593                 ciErr |= clSetKernelArg(ckPathTraceKernel, narg++, sizeof(d_rng_state), (void*)&d_rng_state);
594
595 #define KERNEL_TEX(type, ttype, name) \
596         ciErr |= set_kernel_arg_mem(ckPathTraceKernel, &narg, #name);
597 #include "kernel_textures.h"
598
599                 ciErr |= clSetKernelArg(ckPathTraceKernel, narg++, sizeof(d_sample), (void*)&d_sample);
600                 ciErr |= clSetKernelArg(ckPathTraceKernel, narg++, sizeof(d_x), (void*)&d_x);
601                 ciErr |= clSetKernelArg(ckPathTraceKernel, narg++, sizeof(d_y), (void*)&d_y);
602                 ciErr |= clSetKernelArg(ckPathTraceKernel, narg++, sizeof(d_w), (void*)&d_w);
603                 ciErr |= clSetKernelArg(ckPathTraceKernel, narg++, sizeof(d_h), (void*)&d_h);
604                 ciErr |= clSetKernelArg(ckPathTraceKernel, narg++, sizeof(d_offset), (void*)&d_offset);
605                 ciErr |= clSetKernelArg(ckPathTraceKernel, narg++, sizeof(d_stride), (void*)&d_stride);
606
607                 opencl_assert(ciErr);
608
609                 size_t workgroup_size;
610
611                 clGetKernelWorkGroupInfo(ckPathTraceKernel, cdDevice,
612                         CL_KERNEL_WORK_GROUP_SIZE, sizeof(size_t), &workgroup_size, NULL);
613         
614                 workgroup_size = max(sqrt((double)workgroup_size), 1.0);
615
616                 size_t local_size[2] = {workgroup_size, workgroup_size};
617                 size_t global_size[2] = {global_size_round_up(local_size[0], d_w), global_size_round_up(local_size[1], d_h)};
618
619                 /* run kernel */
620                 ciErr = clEnqueueNDRangeKernel(cqCommandQueue, ckPathTraceKernel, 2, NULL, global_size, local_size, 0, NULL, NULL);
621                 opencl_assert(ciErr);
622                 opencl_assert(clFinish(cqCommandQueue));
623         }
624
625         cl_int set_kernel_arg_mem(cl_kernel kernel, int *narg, const char *name)
626         {
627                 cl_mem ptr;
628                 cl_int err = 0;
629
630                 if(mem_map.find(name) != mem_map.end()) {
631                         device_memory *mem = mem_map[name];
632                 
633                         ptr = CL_MEM_PTR(mem->device_pointer);
634                 }
635                 else {
636                         /* work around NULL not working, even though the spec says otherwise */
637                         ptr = CL_MEM_PTR(null_mem);
638                 }
639                 
640                 err |= clSetKernelArg(kernel, (*narg)++, sizeof(ptr), (void*)&ptr);
641                 opencl_assert(err);
642
643                 return err;
644         }
645
646         void tonemap(DeviceTask& task, device_ptr buffer, device_ptr rgba)
647         {
648                 /* cast arguments to cl types */
649                 cl_mem d_data = CL_MEM_PTR(const_mem_map["__data"]->device_pointer);
650                 cl_mem d_rgba = CL_MEM_PTR(rgba);
651                 cl_mem d_buffer = CL_MEM_PTR(buffer);
652                 cl_int d_x = task.x;
653                 cl_int d_y = task.y;
654                 cl_int d_w = task.w;
655                 cl_int d_h = task.h;
656                 cl_int d_sample = task.sample;
657                 cl_int d_resolution = task.resolution;
658                 cl_int d_offset = task.offset;
659                 cl_int d_stride = task.stride;
660
661                 /* sample arguments */
662                 int narg = 0;
663                 ciErr = 0;
664
665                 ciErr |= clSetKernelArg(ckFilmConvertKernel, narg++, sizeof(d_data), (void*)&d_data);
666                 ciErr |= clSetKernelArg(ckFilmConvertKernel, narg++, sizeof(d_rgba), (void*)&d_rgba);
667                 ciErr |= clSetKernelArg(ckFilmConvertKernel, narg++, sizeof(d_buffer), (void*)&d_buffer);
668
669 #define KERNEL_TEX(type, ttype, name) \
670         ciErr |= set_kernel_arg_mem(ckFilmConvertKernel, &narg, #name);
671 #include "kernel_textures.h"
672
673                 ciErr |= clSetKernelArg(ckFilmConvertKernel, narg++, sizeof(d_sample), (void*)&d_sample);
674                 ciErr |= clSetKernelArg(ckFilmConvertKernel, narg++, sizeof(d_resolution), (void*)&d_resolution);
675                 ciErr |= clSetKernelArg(ckFilmConvertKernel, narg++, sizeof(d_x), (void*)&d_x);
676                 ciErr |= clSetKernelArg(ckFilmConvertKernel, narg++, sizeof(d_y), (void*)&d_y);
677                 ciErr |= clSetKernelArg(ckFilmConvertKernel, narg++, sizeof(d_w), (void*)&d_w);
678                 ciErr |= clSetKernelArg(ckFilmConvertKernel, narg++, sizeof(d_h), (void*)&d_h);
679                 ciErr |= clSetKernelArg(ckFilmConvertKernel, narg++, sizeof(d_offset), (void*)&d_offset);
680                 ciErr |= clSetKernelArg(ckFilmConvertKernel, narg++, sizeof(d_stride), (void*)&d_stride);
681
682                 opencl_assert(ciErr);
683
684                 size_t workgroup_size;
685
686                 clGetKernelWorkGroupInfo(ckFilmConvertKernel, cdDevice,
687                         CL_KERNEL_WORK_GROUP_SIZE, sizeof(size_t), &workgroup_size, NULL);
688         
689                 workgroup_size = max(sqrt((double)workgroup_size), 1.0);
690
691                 size_t local_size[2] = {workgroup_size, workgroup_size};
692                 size_t global_size[2] = {global_size_round_up(local_size[0], d_w), global_size_round_up(local_size[1], d_h)};
693
694                 /* run kernel */
695                 ciErr = clEnqueueNDRangeKernel(cqCommandQueue, ckFilmConvertKernel, 2, NULL, global_size, local_size, 0, NULL, NULL);
696                 opencl_assert(ciErr);
697                 opencl_assert(clFinish(cqCommandQueue));
698         }
699
700         void thread_run(DeviceTask *task)
701         {
702                 if(task->type == DeviceTask::TONEMAP) {
703                         tonemap(*task, task->buffer, task->rgba);
704                 }
705                 else if(task->type == DeviceTask::PATH_TRACE) {
706                         RenderTile tile;
707                         
708                         /* keep rendering tiles until done */
709                         while(task->acquire_tile(this, tile)) {
710                                 int start_sample = tile.start_sample;
711                                 int end_sample = tile.start_sample + tile.num_samples;
712
713                                 for(int sample = start_sample; sample < end_sample; sample++) {
714                                         if (task->get_cancel()) {
715                                                 if(task->need_finish_queue == false)
716                                                         break;
717                                         }
718
719                                         path_trace(tile, sample);
720
721                                         tile.sample = sample + 1;
722
723                                         task->update_progress(tile);
724                                 }
725
726                                 task->release_tile(tile);
727                         }
728                 }
729         }
730
731         class OpenCLDeviceTask : public DeviceTask {
732         public:
733                 OpenCLDeviceTask(OpenCLDevice *device, DeviceTask& task)
734                 : DeviceTask(task)
735                 {
736                         run = function_bind(&OpenCLDevice::thread_run, device, this);
737                 }
738         };
739
740         void task_add(DeviceTask& task)
741         {
742                 task_pool.push(new OpenCLDeviceTask(this, task));
743         }
744
745         void task_wait()
746         {
747                 task_pool.wait_work();
748         }
749
750         void task_cancel()
751         {
752                 task_pool.cancel();
753         }
754 };
755
756 Device *device_opencl_create(DeviceInfo& info, Stats &stats, bool background)
757 {
758         return new OpenCLDevice(info, stats, background);
759 }
760
761 void device_opencl_info(vector<DeviceInfo>& devices)
762 {
763         vector<cl_device_id> device_ids;
764         cl_uint num_devices = 0;
765         vector<cl_platform_id> platform_ids;
766         cl_uint num_platforms = 0;
767
768         /* get devices */
769         if(clGetPlatformIDs(0, NULL, &num_platforms) != CL_SUCCESS || num_platforms == 0)
770                 return;
771         
772         platform_ids.resize(num_platforms);
773
774         if(clGetPlatformIDs(num_platforms, &platform_ids[0], NULL) != CL_SUCCESS)
775                 return;
776
777         if(clGetDeviceIDs(platform_ids[0], opencl_device_type(), 0, NULL, &num_devices) != CL_SUCCESS || num_devices == 0)
778                 return;
779         
780         device_ids.resize(num_devices);
781
782         if(clGetDeviceIDs(platform_ids[0], opencl_device_type(), num_devices, &device_ids[0], NULL) != CL_SUCCESS)
783                 return;
784         
785         /* add devices */
786         for(int num = 0; num < num_devices; num++) {
787                 cl_device_id device_id = device_ids[num];
788                 char name[1024] = "\0";
789
790                 if(clGetDeviceInfo(device_id, CL_DEVICE_NAME, sizeof(name), &name, NULL) != CL_SUCCESS)
791                         continue;
792
793                 DeviceInfo info;
794
795                 info.type = DEVICE_OPENCL;
796                 info.description = string(name);
797                 info.id = string_printf("OPENCL_%d", num);
798                 info.num = num;
799                 /* we don't know if it's used for display, but assume it is */
800                 info.display_device = true;
801                 info.advanced_shading = false;
802                 info.pack_images = true;
803
804                 devices.push_back(info);
805         }
806 }
807
808 CCL_NAMESPACE_END
809
810 #endif /* WITH_OPENCL */
811