5fd67a3ba5
CUDA cards with shader model >= 2 for now (GTX 4xx, 5xx, ..). The CUDA compiler can't handle the increased kernel size currently.
649 lines
19 KiB
C++
649 lines
19 KiB
C++
/*
|
|
* Copyright 2011, Blender Foundation.
|
|
*
|
|
* This program is free software; you can redistribute it and/or
|
|
* modify it under the terms of the GNU General Public License
|
|
* as published by the Free Software Foundation; either version 2
|
|
* of the License, or (at your option) any later version.
|
|
*
|
|
* This program is distributed in the hope that it will be useful,
|
|
* but WITHOUT ANY WARRANTY; without even the implied warranty of
|
|
* MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
|
|
* GNU General Public License for more details.
|
|
*
|
|
* You should have received a copy of the GNU General Public License
|
|
* along with this program; if not, write to the Free Software Foundation,
|
|
* Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
|
|
*/
|
|
|
|
#ifdef WITH_OPENCL
|
|
|
|
#include <stdio.h>
|
|
#include <stdlib.h>
|
|
#include <string.h>
|
|
|
|
#include "device.h"
|
|
#include "device_intern.h"
|
|
|
|
#include "util_map.h"
|
|
#include "util_math.h"
|
|
#include "util_md5.h"
|
|
#include "util_opencl.h"
|
|
#include "util_opengl.h"
|
|
#include "util_path.h"
|
|
#include "util_time.h"
|
|
|
|
CCL_NAMESPACE_BEGIN
|
|
|
|
#define CL_MEM_PTR(p) ((cl_mem)(unsigned long)(p))
|
|
|
|
class OpenCLDevice : public Device
|
|
{
|
|
public:
|
|
cl_context cxContext;
|
|
cl_command_queue cqCommandQueue;
|
|
cl_platform_id cpPlatform;
|
|
cl_device_id cdDevice;
|
|
cl_program cpProgram;
|
|
cl_kernel ckPathTraceKernel;
|
|
cl_kernel ckFilmConvertKernel;
|
|
cl_int ciErr;
|
|
map<string, device_vector<uchar>*> const_mem_map;
|
|
map<string, device_memory*> mem_map;
|
|
device_ptr null_mem;
|
|
bool device_initialized;
|
|
|
|
const char *opencl_error_string(cl_int err)
|
|
{
|
|
switch (err) {
|
|
case CL_SUCCESS: return "Success!";
|
|
case CL_DEVICE_NOT_FOUND: return "Device not found.";
|
|
case CL_DEVICE_NOT_AVAILABLE: return "Device not available";
|
|
case CL_COMPILER_NOT_AVAILABLE: return "Compiler not available";
|
|
case CL_MEM_OBJECT_ALLOCATION_FAILURE: return "Memory object allocation failure";
|
|
case CL_OUT_OF_RESOURCES: return "Out of resources";
|
|
case CL_OUT_OF_HOST_MEMORY: return "Out of host memory";
|
|
case CL_PROFILING_INFO_NOT_AVAILABLE: return "Profiling information not available";
|
|
case CL_MEM_COPY_OVERLAP: return "Memory copy overlap";
|
|
case CL_IMAGE_FORMAT_MISMATCH: return "Image format mismatch";
|
|
case CL_IMAGE_FORMAT_NOT_SUPPORTED: return "Image format not supported";
|
|
case CL_BUILD_PROGRAM_FAILURE: return "Program build failure";
|
|
case CL_MAP_FAILURE: return "Map failure";
|
|
case CL_INVALID_VALUE: return "Invalid value";
|
|
case CL_INVALID_DEVICE_TYPE: return "Invalid device type";
|
|
case CL_INVALID_PLATFORM: return "Invalid platform";
|
|
case CL_INVALID_DEVICE: return "Invalid device";
|
|
case CL_INVALID_CONTEXT: return "Invalid context";
|
|
case CL_INVALID_QUEUE_PROPERTIES: return "Invalid queue properties";
|
|
case CL_INVALID_COMMAND_QUEUE: return "Invalid command queue";
|
|
case CL_INVALID_HOST_PTR: return "Invalid host pointer";
|
|
case CL_INVALID_MEM_OBJECT: return "Invalid memory object";
|
|
case CL_INVALID_IMAGE_FORMAT_DESCRIPTOR: return "Invalid image format descriptor";
|
|
case CL_INVALID_IMAGE_SIZE: return "Invalid image size";
|
|
case CL_INVALID_SAMPLER: return "Invalid sampler";
|
|
case CL_INVALID_BINARY: return "Invalid binary";
|
|
case CL_INVALID_BUILD_OPTIONS: return "Invalid build options";
|
|
case CL_INVALID_PROGRAM: return "Invalid program";
|
|
case CL_INVALID_PROGRAM_EXECUTABLE: return "Invalid program executable";
|
|
case CL_INVALID_KERNEL_NAME: return "Invalid kernel name";
|
|
case CL_INVALID_KERNEL_DEFINITION: return "Invalid kernel definition";
|
|
case CL_INVALID_KERNEL: return "Invalid kernel";
|
|
case CL_INVALID_ARG_INDEX: return "Invalid argument index";
|
|
case CL_INVALID_ARG_VALUE: return "Invalid argument value";
|
|
case CL_INVALID_ARG_SIZE: return "Invalid argument size";
|
|
case CL_INVALID_KERNEL_ARGS: return "Invalid kernel arguments";
|
|
case CL_INVALID_WORK_DIMENSION: return "Invalid work dimension";
|
|
case CL_INVALID_WORK_GROUP_SIZE: return "Invalid work group size";
|
|
case CL_INVALID_WORK_ITEM_SIZE: return "Invalid work item size";
|
|
case CL_INVALID_GLOBAL_OFFSET: return "Invalid global offset";
|
|
case CL_INVALID_EVENT_WAIT_LIST: return "Invalid event wait list";
|
|
case CL_INVALID_EVENT: return "Invalid event";
|
|
case CL_INVALID_OPERATION: return "Invalid operation";
|
|
case CL_INVALID_GL_OBJECT: return "Invalid OpenGL object";
|
|
case CL_INVALID_BUFFER_SIZE: return "Invalid buffer size";
|
|
case CL_INVALID_MIP_LEVEL: return "Invalid mip-map level";
|
|
default: return "Unknown";
|
|
}
|
|
}
|
|
|
|
bool opencl_error(cl_int err)
|
|
{
|
|
if(err != CL_SUCCESS) {
|
|
fprintf(stderr, "OpenCL error (%d): %s\n", err, opencl_error_string(err));
|
|
return true;
|
|
}
|
|
|
|
return false;
|
|
}
|
|
|
|
void opencl_assert(cl_int err)
|
|
{
|
|
if(err != CL_SUCCESS) {
|
|
fprintf(stderr, "OpenCL error (%d): %s\n", err, opencl_error_string(err));
|
|
#ifndef NDEBUG
|
|
abort();
|
|
#endif
|
|
}
|
|
}
|
|
|
|
OpenCLDevice(bool background_)
|
|
{
|
|
background = background_;
|
|
cpPlatform = NULL;
|
|
cxContext = NULL;
|
|
cqCommandQueue = NULL;
|
|
cpProgram = NULL;
|
|
ckPathTraceKernel = NULL;
|
|
ckFilmConvertKernel = NULL;
|
|
null_mem = 0;
|
|
device_initialized = false;
|
|
|
|
vector<cl_platform_id> platform_ids;
|
|
cl_uint num_platforms;
|
|
|
|
/* setup device */
|
|
ciErr = clGetPlatformIDs(0, NULL, &num_platforms);
|
|
if(opencl_error(ciErr))
|
|
return;
|
|
|
|
if(num_platforms == 0) {
|
|
fprintf(stderr, "OpenCL: no platforms found.\n");
|
|
return;
|
|
}
|
|
|
|
platform_ids.resize(num_platforms);
|
|
ciErr = clGetPlatformIDs(num_platforms, &platform_ids[0], NULL);
|
|
if(opencl_error(ciErr))
|
|
return;
|
|
|
|
cpPlatform = platform_ids[0]; /* todo: pick specified platform && device */
|
|
|
|
ciErr = clGetDeviceIDs(cpPlatform, CL_DEVICE_TYPE_GPU|CL_DEVICE_TYPE_ACCELERATOR, 1, &cdDevice, NULL);
|
|
if(opencl_error(ciErr))
|
|
return;
|
|
|
|
cxContext = clCreateContext(0, 1, &cdDevice, NULL, NULL, &ciErr);
|
|
if(opencl_error(ciErr))
|
|
return;
|
|
|
|
cqCommandQueue = clCreateCommandQueue(cxContext, cdDevice, 0, &ciErr);
|
|
if(opencl_error(ciErr))
|
|
return;
|
|
|
|
null_mem = (device_ptr)clCreateBuffer(cxContext, CL_MEM_READ_ONLY, 1, NULL, &ciErr);
|
|
device_initialized = true;
|
|
}
|
|
|
|
bool opencl_version_check()
|
|
{
|
|
char version[256];
|
|
int major, minor, req_major = 1, req_minor = 1;
|
|
|
|
clGetPlatformInfo(cpPlatform, CL_PLATFORM_VERSION, sizeof(version), &version, NULL);
|
|
|
|
if(sscanf(version, "OpenCL %d.%d", &major, &minor) < 2) {
|
|
fprintf(stderr, "OpenCL: failed to parse platform version string (%s).", version);
|
|
return false;
|
|
}
|
|
|
|
if(!((major == req_major && minor >= req_minor) || (major > req_major))) {
|
|
fprintf(stderr, "OpenCL: platform version 1.1 or later required, found %d.%d\n", major, minor);
|
|
return false;
|
|
}
|
|
|
|
clGetDeviceInfo(cdDevice, CL_DEVICE_OPENCL_C_VERSION, sizeof(version), &version, NULL);
|
|
|
|
if(sscanf(version, "OpenCL C %d.%d", &major, &minor) < 2) {
|
|
fprintf(stderr, "OpenCL: failed to parse OpenCL C version string (%s).", version);
|
|
return false;
|
|
}
|
|
|
|
if(!((major == req_major && minor >= req_minor) || (major > req_major))) {
|
|
fprintf(stderr, "OpenCL: C version 1.1 or later required, found %d.%d\n", major, minor);
|
|
return false;
|
|
}
|
|
|
|
/* we don't check CL_DEVICE_VERSION since for e.g. nvidia sm 1.3 cards this is
|
|
1.0 even if the language features are there, just limited shared memory */
|
|
|
|
return true;
|
|
}
|
|
|
|
bool load_binary(const string& kernel_path, const string& clbin)
|
|
{
|
|
/* read binary into memory */
|
|
vector<uint8_t> binary;
|
|
|
|
if(!path_read_binary(clbin, binary)) {
|
|
fprintf(stderr, "OpenCL failed to read cached binary %s.\n", clbin.c_str());
|
|
return false;
|
|
}
|
|
|
|
/* create program */
|
|
cl_int status;
|
|
size_t size = binary.size();
|
|
const uint8_t *bytes = &binary[0];
|
|
|
|
cpProgram = clCreateProgramWithBinary(cxContext, 1, &cdDevice,
|
|
&size, &bytes, &status, &ciErr);
|
|
|
|
if(opencl_error(status) || opencl_error(ciErr)) {
|
|
fprintf(stderr, "OpenCL failed create program from cached binary %s.\n", clbin.c_str());
|
|
return false;
|
|
}
|
|
|
|
if(!build_kernel(kernel_path))
|
|
return false;
|
|
|
|
return true;
|
|
}
|
|
|
|
bool save_binary(const string& clbin)
|
|
{
|
|
size_t size = 0;
|
|
clGetProgramInfo(cpProgram, CL_PROGRAM_BINARY_SIZES, sizeof(size_t), &size, NULL);
|
|
|
|
if(!size)
|
|
return false;
|
|
|
|
vector<uint8_t> binary(size);
|
|
uint8_t *bytes = &binary[0];
|
|
|
|
clGetProgramInfo(cpProgram, CL_PROGRAM_BINARIES, sizeof(uint8_t*), &bytes, NULL);
|
|
|
|
if(!path_write_binary(clbin, binary)) {
|
|
fprintf(stderr, "OpenCL failed to write cached binary %s.\n", clbin.c_str());
|
|
return false;
|
|
}
|
|
|
|
return true;
|
|
}
|
|
|
|
bool build_kernel(const string& kernel_path)
|
|
{
|
|
string build_options = "";
|
|
|
|
build_options += "-I " + kernel_path + ""; /* todo: escape path */
|
|
build_options += " -cl-fast-relaxed-math ";
|
|
|
|
ciErr = clBuildProgram(cpProgram, 0, NULL, build_options.c_str(), NULL, NULL);
|
|
|
|
if(ciErr != CL_SUCCESS) {
|
|
/* show build errors */
|
|
char *build_log;
|
|
size_t ret_val_size;
|
|
|
|
clGetProgramBuildInfo(cpProgram, cdDevice, CL_PROGRAM_BUILD_LOG, 0, NULL, &ret_val_size);
|
|
|
|
build_log = new char[ret_val_size+1];
|
|
clGetProgramBuildInfo(cpProgram, cdDevice, CL_PROGRAM_BUILD_LOG, ret_val_size, build_log, NULL);
|
|
|
|
build_log[ret_val_size] = '\0';
|
|
fprintf(stderr, "OpenCL build failed:\n %s\n", build_log);
|
|
|
|
delete[] build_log;
|
|
|
|
return false;
|
|
}
|
|
|
|
return true;
|
|
}
|
|
|
|
bool compile_kernel(const string& kernel_path, const string& kernel_md5)
|
|
{
|
|
/* we compile kernels consisting of many files. unfortunately opencl
|
|
kernel caches do not seem to recognize changes in included files.
|
|
so we force recompile on changes by adding the md5 hash of all files */
|
|
string source = "#include \"kernel.cl\" // " + kernel_md5 + "\n";
|
|
size_t source_len = source.size();
|
|
const char *source_str = source.c_str();
|
|
|
|
cpProgram = clCreateProgramWithSource(cxContext, 1, &source_str, &source_len, &ciErr);
|
|
|
|
if(opencl_error(ciErr))
|
|
return false;
|
|
|
|
double starttime = time_dt();
|
|
printf("Compiling OpenCL kernel ...\n");
|
|
|
|
if(!build_kernel(kernel_path))
|
|
return false;
|
|
|
|
printf("Kernel compilation finished in %.2lfs.\n", time_dt() - starttime);
|
|
|
|
return true;
|
|
}
|
|
|
|
string device_md5_hash()
|
|
{
|
|
MD5Hash md5;
|
|
char version[256], driver[256], name[256], vendor[256];
|
|
|
|
clGetPlatformInfo(cpPlatform, CL_PLATFORM_VENDOR, sizeof(vendor), &vendor, NULL);
|
|
clGetDeviceInfo(cdDevice, CL_DEVICE_VERSION, sizeof(version), &version, NULL);
|
|
clGetDeviceInfo(cdDevice, CL_DEVICE_NAME, sizeof(name), &name, NULL);
|
|
clGetDeviceInfo(cdDevice, CL_DRIVER_VERSION, sizeof(driver), &driver, NULL);
|
|
|
|
md5.append((uint8_t*)vendor, strlen(vendor));
|
|
md5.append((uint8_t*)version, strlen(version));
|
|
md5.append((uint8_t*)name, strlen(name));
|
|
md5.append((uint8_t*)driver, strlen(driver));
|
|
|
|
return md5.get_hex();
|
|
}
|
|
|
|
bool load_kernels()
|
|
{
|
|
/* verify if device was initialized */
|
|
if(!device_initialized) {
|
|
fprintf(stderr, "OpenCL: failed to initialize device.\n");
|
|
return false;
|
|
}
|
|
|
|
/* verify we have right opencl version */
|
|
if(!opencl_version_check())
|
|
return false;
|
|
|
|
/* md5 hash to detect changes */
|
|
string kernel_path = path_get("kernel");
|
|
string kernel_md5 = path_files_md5_hash(kernel_path);
|
|
string device_md5 = device_md5_hash();
|
|
|
|
/* try to use cache binary */
|
|
string clbin = string_printf("cycles_kernel_%s_%s.clbin", device_md5.c_str(), kernel_md5.c_str());;
|
|
clbin = path_user_get(path_join("cache", clbin));
|
|
|
|
if(path_exists(clbin)) {
|
|
/* if exists already, try use it */
|
|
if(!load_binary(kernel_path, clbin))
|
|
return false;
|
|
}
|
|
else {
|
|
/* compile kernel */
|
|
if(!compile_kernel(kernel_path, kernel_md5))
|
|
return false;
|
|
|
|
/* save binary for reuse */
|
|
save_binary(clbin);
|
|
}
|
|
|
|
/* find kernels */
|
|
ckPathTraceKernel = clCreateKernel(cpProgram, "kernel_ocl_path_trace", &ciErr);
|
|
if(opencl_error(ciErr))
|
|
return false;
|
|
|
|
ckFilmConvertKernel = clCreateKernel(cpProgram, "kernel_ocl_tonemap", &ciErr);
|
|
if(opencl_error(ciErr))
|
|
return false;
|
|
|
|
return true;
|
|
}
|
|
|
|
~OpenCLDevice()
|
|
{
|
|
if(null_mem)
|
|
clReleaseMemObject(CL_MEM_PTR(null_mem));
|
|
|
|
map<string, device_vector<uchar>*>::iterator mt;
|
|
for(mt = const_mem_map.begin(); mt != const_mem_map.end(); mt++) {
|
|
mem_free(*(mt->second));
|
|
delete mt->second;
|
|
}
|
|
|
|
if(ckPathTraceKernel)
|
|
clReleaseKernel(ckPathTraceKernel);
|
|
if(ckFilmConvertKernel)
|
|
clReleaseKernel(ckFilmConvertKernel);
|
|
if(cpProgram)
|
|
clReleaseProgram(cpProgram);
|
|
if(cqCommandQueue)
|
|
clReleaseCommandQueue(cqCommandQueue);
|
|
if(cxContext)
|
|
clReleaseContext(cxContext);
|
|
}
|
|
|
|
bool support_full_kernel()
|
|
{
|
|
return false;
|
|
}
|
|
|
|
string description()
|
|
{
|
|
char name[1024];
|
|
|
|
clGetDeviceInfo(cdDevice, CL_DEVICE_NAME, sizeof(name), &name, NULL);
|
|
|
|
return string("OpenCL ") + name;
|
|
}
|
|
|
|
void mem_alloc(device_memory& mem, MemoryType type)
|
|
{
|
|
size_t size = mem.memory_size();
|
|
|
|
if(type == MEM_READ_ONLY)
|
|
mem.device_pointer = (device_ptr)clCreateBuffer(cxContext, CL_MEM_READ_ONLY, size, NULL, &ciErr);
|
|
else if(type == MEM_WRITE_ONLY)
|
|
mem.device_pointer = (device_ptr)clCreateBuffer(cxContext, CL_MEM_WRITE_ONLY, size, NULL, &ciErr);
|
|
else
|
|
mem.device_pointer = (device_ptr)clCreateBuffer(cxContext, CL_MEM_READ_WRITE, size, NULL, &ciErr);
|
|
|
|
opencl_assert(ciErr);
|
|
}
|
|
|
|
void mem_copy_to(device_memory& mem)
|
|
{
|
|
/* this is blocking */
|
|
size_t size = mem.memory_size();
|
|
ciErr = clEnqueueWriteBuffer(cqCommandQueue, CL_MEM_PTR(mem.device_pointer), CL_TRUE, 0, size, (void*)mem.data_pointer, 0, NULL, NULL);
|
|
opencl_assert(ciErr);
|
|
}
|
|
|
|
void mem_copy_from(device_memory& mem, size_t offset, size_t size)
|
|
{
|
|
ciErr = clEnqueueReadBuffer(cqCommandQueue, CL_MEM_PTR(mem.device_pointer), CL_TRUE, offset, size, (uchar*)mem.data_pointer + offset, 0, NULL, NULL);
|
|
opencl_assert(ciErr);
|
|
}
|
|
|
|
void mem_zero(device_memory& mem)
|
|
{
|
|
if(mem.device_pointer) {
|
|
memset((void*)mem.data_pointer, 0, mem.memory_size());
|
|
mem_copy_to(mem);
|
|
}
|
|
}
|
|
|
|
void mem_free(device_memory& mem)
|
|
{
|
|
if(mem.device_pointer) {
|
|
ciErr = clReleaseMemObject(CL_MEM_PTR(mem.device_pointer));
|
|
mem.device_pointer = 0;
|
|
opencl_assert(ciErr);
|
|
}
|
|
}
|
|
|
|
void const_copy_to(const char *name, void *host, size_t size)
|
|
{
|
|
if(const_mem_map.find(name) == const_mem_map.end()) {
|
|
device_vector<uchar> *data = new device_vector<uchar>();
|
|
data->copy((uchar*)host, size);
|
|
|
|
mem_alloc(*data, MEM_READ_ONLY);
|
|
const_mem_map[name] = data;
|
|
}
|
|
else {
|
|
device_vector<uchar> *data = const_mem_map[name];
|
|
data->copy((uchar*)host, size);
|
|
}
|
|
|
|
mem_copy_to(*const_mem_map[name]);
|
|
}
|
|
|
|
void tex_alloc(const char *name, device_memory& mem, bool interpolation, bool periodic)
|
|
{
|
|
mem_alloc(mem, MEM_READ_ONLY);
|
|
mem_copy_to(mem);
|
|
mem_map[name] = &mem;
|
|
}
|
|
|
|
void tex_free(device_memory& mem)
|
|
{
|
|
if(mem.data_pointer)
|
|
mem_free(mem);
|
|
}
|
|
|
|
size_t global_size_round_up(int group_size, int global_size)
|
|
{
|
|
int r = global_size % group_size;
|
|
return global_size + ((r == 0)? 0: group_size - r);
|
|
}
|
|
|
|
void path_trace(DeviceTask& task)
|
|
{
|
|
/* cast arguments to cl types */
|
|
cl_mem d_data = CL_MEM_PTR(const_mem_map["__data"]->device_pointer);
|
|
cl_mem d_buffer = CL_MEM_PTR(task.buffer);
|
|
cl_mem d_rng_state = CL_MEM_PTR(task.rng_state);
|
|
cl_int d_x = task.x;
|
|
cl_int d_y = task.y;
|
|
cl_int d_w = task.w;
|
|
cl_int d_h = task.h;
|
|
cl_int d_sample = task.sample;
|
|
|
|
/* sample arguments */
|
|
int narg = 0;
|
|
ciErr = 0;
|
|
|
|
ciErr |= clSetKernelArg(ckPathTraceKernel, narg++, sizeof(d_data), (void*)&d_data);
|
|
ciErr |= clSetKernelArg(ckPathTraceKernel, narg++, sizeof(d_buffer), (void*)&d_buffer);
|
|
ciErr |= clSetKernelArg(ckPathTraceKernel, narg++, sizeof(d_rng_state), (void*)&d_rng_state);
|
|
|
|
#define KERNEL_TEX(type, ttype, name) \
|
|
ciErr |= set_kernel_arg_mem(ckPathTraceKernel, &narg, #name);
|
|
#include "kernel_textures.h"
|
|
|
|
ciErr |= clSetKernelArg(ckPathTraceKernel, narg++, sizeof(d_sample), (void*)&d_sample);
|
|
ciErr |= clSetKernelArg(ckPathTraceKernel, narg++, sizeof(d_x), (void*)&d_x);
|
|
ciErr |= clSetKernelArg(ckPathTraceKernel, narg++, sizeof(d_y), (void*)&d_y);
|
|
ciErr |= clSetKernelArg(ckPathTraceKernel, narg++, sizeof(d_w), (void*)&d_w);
|
|
ciErr |= clSetKernelArg(ckPathTraceKernel, narg++, sizeof(d_h), (void*)&d_h);
|
|
|
|
opencl_assert(ciErr);
|
|
|
|
size_t workgroup_size;
|
|
|
|
clGetKernelWorkGroupInfo(ckPathTraceKernel, cdDevice,
|
|
CL_KERNEL_WORK_GROUP_SIZE, sizeof(size_t), &workgroup_size, NULL);
|
|
|
|
workgroup_size = max(sqrt((double)workgroup_size), 1.0);
|
|
|
|
size_t local_size[2] = {workgroup_size, workgroup_size};
|
|
size_t global_size[2] = {global_size_round_up(local_size[0], d_w), global_size_round_up(local_size[1], d_h)};
|
|
|
|
/* run kernel */
|
|
ciErr = clEnqueueNDRangeKernel(cqCommandQueue, ckPathTraceKernel, 2, NULL, global_size, local_size, 0, NULL, NULL);
|
|
opencl_assert(ciErr);
|
|
opencl_assert(clFinish(cqCommandQueue));
|
|
}
|
|
|
|
cl_int set_kernel_arg_mem(cl_kernel kernel, int *narg, const char *name)
|
|
{
|
|
cl_mem ptr;
|
|
cl_int size, err = 0;
|
|
|
|
if(mem_map.find(name) != mem_map.end()) {
|
|
device_memory *mem = mem_map[name];
|
|
|
|
ptr = CL_MEM_PTR(mem->device_pointer);
|
|
size = mem->data_width;
|
|
}
|
|
else {
|
|
/* work around NULL not working, even though the spec says otherwise */
|
|
ptr = CL_MEM_PTR(null_mem);
|
|
size = 1;
|
|
}
|
|
|
|
err |= clSetKernelArg(kernel, (*narg)++, sizeof(ptr), (void*)&ptr);
|
|
opencl_assert(err);
|
|
err |= clSetKernelArg(kernel, (*narg)++, sizeof(size), (void*)&size);
|
|
opencl_assert(err);
|
|
|
|
return err;
|
|
}
|
|
|
|
void tonemap(DeviceTask& task)
|
|
{
|
|
/* cast arguments to cl types */
|
|
cl_mem d_data = CL_MEM_PTR(const_mem_map["__data"]->device_pointer);
|
|
cl_mem d_rgba = CL_MEM_PTR(task.rgba);
|
|
cl_mem d_buffer = CL_MEM_PTR(task.buffer);
|
|
cl_int d_x = task.x;
|
|
cl_int d_y = task.y;
|
|
cl_int d_w = task.w;
|
|
cl_int d_h = task.h;
|
|
cl_int d_sample = task.sample;
|
|
cl_int d_resolution = task.resolution;
|
|
|
|
/* sample arguments */
|
|
int narg = 0;
|
|
ciErr = 0;
|
|
|
|
ciErr |= clSetKernelArg(ckFilmConvertKernel, narg++, sizeof(d_data), (void*)&d_data);
|
|
ciErr |= clSetKernelArg(ckFilmConvertKernel, narg++, sizeof(d_rgba), (void*)&d_rgba);
|
|
ciErr |= clSetKernelArg(ckFilmConvertKernel, narg++, sizeof(d_buffer), (void*)&d_buffer);
|
|
|
|
#define KERNEL_TEX(type, ttype, name) \
|
|
ciErr |= set_kernel_arg_mem(ckFilmConvertKernel, &narg, #name);
|
|
#include "kernel_textures.h"
|
|
|
|
ciErr |= clSetKernelArg(ckFilmConvertKernel, narg++, sizeof(d_sample), (void*)&d_sample);
|
|
ciErr |= clSetKernelArg(ckFilmConvertKernel, narg++, sizeof(d_resolution), (void*)&d_resolution);
|
|
ciErr |= clSetKernelArg(ckFilmConvertKernel, narg++, sizeof(d_x), (void*)&d_x);
|
|
ciErr |= clSetKernelArg(ckFilmConvertKernel, narg++, sizeof(d_y), (void*)&d_y);
|
|
ciErr |= clSetKernelArg(ckFilmConvertKernel, narg++, sizeof(d_w), (void*)&d_w);
|
|
ciErr |= clSetKernelArg(ckFilmConvertKernel, narg++, sizeof(d_h), (void*)&d_h);
|
|
|
|
opencl_assert(ciErr);
|
|
|
|
size_t workgroup_size;
|
|
|
|
clGetKernelWorkGroupInfo(ckFilmConvertKernel, cdDevice,
|
|
CL_KERNEL_WORK_GROUP_SIZE, sizeof(size_t), &workgroup_size, NULL);
|
|
|
|
workgroup_size = max(sqrt((double)workgroup_size), 1.0);
|
|
|
|
size_t local_size[2] = {workgroup_size, workgroup_size};
|
|
size_t global_size[2] = {global_size_round_up(local_size[0], d_w), global_size_round_up(local_size[1], d_h)};
|
|
|
|
/* run kernel */
|
|
ciErr = clEnqueueNDRangeKernel(cqCommandQueue, ckFilmConvertKernel, 2, NULL, global_size, local_size, 0, NULL, NULL);
|
|
opencl_assert(ciErr);
|
|
opencl_assert(clFinish(cqCommandQueue));
|
|
}
|
|
|
|
void task_add(DeviceTask& task)
|
|
{
|
|
if(task.type == DeviceTask::TONEMAP)
|
|
tonemap(task);
|
|
else if(task.type == DeviceTask::PATH_TRACE)
|
|
path_trace(task);
|
|
}
|
|
|
|
void task_wait()
|
|
{
|
|
}
|
|
|
|
void task_cancel()
|
|
{
|
|
}
|
|
};
|
|
|
|
Device *device_opencl_create(bool background)
|
|
{
|
|
return new OpenCLDevice(background);
|
|
}
|
|
|
|
CCL_NAMESPACE_END
|
|
|
|
#endif /* WITH_OPENCL */
|
|
|