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