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