39 OpenCLCache::Slot::ProgramEntry::ProgramEntry(
const ProgramEntry &rhs)
44 OpenCLCache::Slot::ProgramEntry::~ProgramEntry()
53 OpenCLCache::Slot::Slot(
const Slot &rhs)
58 OpenCLCache::Slot::~Slot()
63 OpenCLCache &OpenCLCache::global_instance()
65 static OpenCLCache instance;
69 cl_context OpenCLCache::get_context(cl_platform_id platform,
73 assert(platform !=
NULL);
75 OpenCLCache &
self = global_instance();
79 pair<CacheMap::iterator, bool> ins =
self.cache.insert(
80 CacheMap::value_type(PlatformDevicePair(platform, device), Slot()));
82 Slot &slot = ins.first->second;
85 if (!slot.context_mutex)
95 if (slot.context ==
NULL) {
101 slot_locker.unlock();
103 cl_int ciErr = clRetainContext(slot.context);
104 assert(ciErr == CL_SUCCESS);
110 cl_program OpenCLCache::get_program(cl_platform_id platform,
115 assert(platform !=
NULL);
117 OpenCLCache &
self = global_instance();
121 pair<CacheMap::iterator, bool> ins =
self.cache.insert(
122 CacheMap::value_type(PlatformDevicePair(platform, device), Slot()));
124 Slot &slot = ins.first->second;
126 pair<Slot::EntryMap::iterator, bool> ins2 = slot.programs.insert(
127 Slot::EntryMap::value_type(key, Slot::ProgramEntry()));
129 Slot::ProgramEntry &entry = ins2.first->second;
142 if (entry.program ==
NULL) {
148 slot_locker.unlock();
150 cl_int ciErr = clRetainProgram(entry.program);
151 assert(ciErr == CL_SUCCESS);
154 return entry.program;
157 void OpenCLCache::store_context(cl_platform_id platform,
162 assert(platform !=
NULL);
163 assert(device !=
NULL);
166 OpenCLCache &
self = global_instance();
169 CacheMap::iterator i =
self.cache.find(PlatformDevicePair(platform, device));
172 Slot &slot = i->second;
175 assert(i !=
self.cache.end());
176 assert(slot.context ==
NULL);
181 slot_locker.unlock();
185 cl_int ciErr = clRetainContext(
context);
186 assert(ciErr == CL_SUCCESS);
190 void OpenCLCache::store_program(cl_platform_id platform,
196 assert(platform !=
NULL);
197 assert(device !=
NULL);
200 OpenCLCache &
self = global_instance();
204 CacheMap::iterator i =
self.cache.find(PlatformDevicePair(platform, device));
205 assert(i !=
self.cache.end());
206 Slot &slot = i->second;
208 Slot::EntryMap::iterator i2 = slot.programs.find(key);
209 assert(i2 != slot.programs.end());
210 Slot::ProgramEntry &entry = i2->second;
212 assert(entry.program ==
NULL);
219 slot_locker.unlock();
224 cl_int ciErr = clRetainProgram(
program);
225 assert(ciErr == CL_SUCCESS);
229 string OpenCLCache::get_kernel_md5()
231 OpenCLCache &
self = global_instance();
234 if (
self.kernel_md5.empty()) {
237 return self.kernel_md5;
240 static string get_program_source(
const string &kernel_file)
242 string source =
"#include \"kernel/kernels/opencl/" + kernel_file +
"\"\n";
252 OpenCLDevice::OpenCLProgram::OpenCLProgram(OpenCLDevice *device,
253 const string &program_name,
254 const string &kernel_file,
255 const string &kernel_build_options,
258 program_name(program_name),
259 kernel_file(kernel_file),
260 kernel_build_options(kernel_build_options),
261 use_stdout(use_stdout)
264 needs_compiling =
true;
268 OpenCLDevice::OpenCLProgram::~OpenCLProgram()
273 void OpenCLDevice::OpenCLProgram::release()
275 for (map<ustring, cl_kernel>::iterator kernel = kernels.begin(); kernel != kernels.end();
277 if (kernel->second) {
278 clReleaseKernel(kernel->second);
279 kernel->second =
NULL;
288 void OpenCLDevice::OpenCLProgram::add_log(
const string &msg,
bool debug)
294 printf(
"%s\n", msg.c_str());
302 void OpenCLDevice::OpenCLProgram::add_error(
const string &msg)
305 fprintf(stderr,
"%s\n", msg.c_str());
307 if (error_msg ==
"") {
313 void OpenCLDevice::OpenCLProgram::add_kernel(ustring name)
315 if (!kernels.count(name)) {
316 kernels[name] =
NULL;
320 bool OpenCLDevice::OpenCLProgram::build_kernel(
const string *debug_src)
322 string build_options;
323 build_options = device->kernel_build_options(debug_src) + kernel_build_options;
325 VLOG(1) <<
"Build options passed to clBuildProgram: '" << build_options <<
"'.";
329 size_t ret_val_size = 0;
331 clGetProgramBuildInfo(
program, device->cdDevice, CL_PROGRAM_BUILD_LOG, 0,
NULL, &ret_val_size);
333 if (ciErr != CL_SUCCESS) {
334 add_error(
string(
"OpenCL build failed with error ") + clewErrorString(ciErr) +
335 ", errors in console.");
338 if (ret_val_size > 1) {
340 clGetProgramBuildInfo(
341 program, device->cdDevice, CL_PROGRAM_BUILD_LOG, ret_val_size, &build_log[0],
NULL);
343 build_log[ret_val_size] =
'\0';
345 if (!(ret_val_size == 2 && build_log[0] ==
'\n')) {
346 add_log(
string(
"OpenCL program ") + program_name +
" build output: " +
string(&build_log[0]),
347 ciErr == CL_SUCCESS);
351 return (ciErr == CL_SUCCESS);
354 bool OpenCLDevice::OpenCLProgram::compile_kernel(
const string *debug_src)
356 string source = get_program_source(kernel_file);
362 size_t source_len = source.size();
363 const char *source_str = source.c_str();
366 program = clCreateProgramWithSource(device->cxContext, 1, &source_str, &source_len, &ciErr);
368 if (ciErr != CL_SUCCESS) {
369 add_error(
string(
"OpenCL program creation failed: ") + clewErrorString(ciErr));
374 add_log(
string(
"Cycles: compiling OpenCL program ") + program_name +
"...",
false);
375 add_log(
string(
"Build flags: ") + kernel_build_options,
true);
377 if (!build_kernel(debug_src))
380 double elapsed =
time_dt() - starttime;
382 string_printf(
"Kernel compilation of %s finished in %.2lfs.", program_name.c_str(), elapsed),
388 static void escape_python_string(
string &
str)
394 static int opencl_compile_process_limit()
403 static const int64_t GB = 1024LL * 1024LL * 1024LL;
404 static const int64_t process_memory = 2 * GB;
405 static const int64_t base_memory = 2 * GB;
407 static const int64_t process_limit = (system_memory - base_memory) / process_memory;
409 return max((
int)process_limit, 1);
412 bool OpenCLDevice::OpenCLProgram::compile_separate(
const string &clbin)
416 args.push_back(
"--background");
417 args.push_back(
"--factory-startup");
418 args.push_back(
"--python-expr");
420 int device_platform_id = device->device_num;
421 string device_name = device->device_name;
422 string platform_name = device->platform_name;
423 string build_options = device->kernel_build_options(
NULL) + kernel_build_options;
424 string kernel_file_escaped = kernel_file;
425 string clbin_escaped = clbin;
427 escape_python_string(device_name);
428 escape_python_string(platform_name);
429 escape_python_string(build_options);
430 escape_python_string(kernel_file_escaped);
431 escape_python_string(clbin_escaped);
434 "import _cycles; _cycles.opencl_compile(r'%d', r'%s', r'%s', r'%s', r'%s', r'%s')",
437 platform_name.c_str(),
438 build_options.c_str(),
439 kernel_file_escaped.c_str(),
440 clbin_escaped.c_str()));
447 const double starttime =
time_dt();
448 add_log(
string(
"Cycles: compiling OpenCL program ") + program_name +
"...",
false);
449 add_log(
string(
"Build flags: ") + kernel_build_options,
true);
451 const double elapsed =
time_dt() - starttime;
460 string_printf(
"Kernel compilation of %s finished in %.2lfs.", program_name.c_str(), elapsed),
463 return load_binary(clbin);
470 int device_platform_id = std::stoi(
parameters[0]);
477 if (clewInit() != CLEW_SUCCESS) {
482 OpenCLInfo::get_usable_devices(&usable_devices);
483 if (device_platform_id >= usable_devices.size()) {
487 OpenCLPlatformDevice &platform_device = usable_devices[device_platform_id];
488 if (platform_device.platform_name != platform_name ||
489 platform_device.device_name != device_name) {
493 cl_platform_id platform = platform_device.platform_id;
494 cl_device_id device = platform_device.device_id;
495 const cl_context_properties context_props[] = {
496 CL_CONTEXT_PLATFORM, (cl_context_properties)platform, 0, 0};
500 if (
err != CL_SUCCESS) {
504 string source = get_program_source(kernel_file);
505 size_t source_len = source.size();
506 const char *source_str = source.c_str();
507 cl_program
program = clCreateProgramWithSource(
context, 1, &source_str, &source_len, &
err);
510 if (
err == CL_SUCCESS) {
513 if (
err == CL_SUCCESS) {
515 clGetProgramInfo(
program, CL_PROGRAM_BINARY_SIZES,
sizeof(
size_t), &
size,
NULL);
531 bool OpenCLDevice::OpenCLProgram::load_binary(
const string &clbin,
const string *debug_src)
537 add_error(
string_printf(
"OpenCL failed to read cached binary %s.", clbin.c_str()));
542 cl_int status, ciErr;
543 size_t size = binary.size();
544 const uint8_t *bytes = &binary[0];
546 program = clCreateProgramWithBinary(
547 device->cxContext, 1, &device->cdDevice, &
size, &bytes, &status, &ciErr);
549 if (status != CL_SUCCESS || ciErr != CL_SUCCESS) {
550 add_error(
string(
"OpenCL failed create program from cached binary ") + clbin +
": " +
551 clewErrorString(status) +
" " + clewErrorString(ciErr));
555 if (!build_kernel(debug_src))
561 bool OpenCLDevice::OpenCLProgram::save_binary(
const string &clbin)
564 clGetProgramInfo(
program, CL_PROGRAM_BINARY_SIZES,
sizeof(
size_t), &
size,
NULL);
580 string device_md5 = device->device_md5_hash(kernel_build_options);
584 ustring cache_key(program_name + device_md5);
585 program = device->load_cached_kernel(cache_key, cache_locker);
587 add_log(
string(
"OpenCL program ") + program_name +
" not found in cache.",
true);
590 string source = get_program_source(kernel_file);
592 string basename =
"cycles_kernel_" + program_name +
"_" + device_md5 +
"_" +
600 add_log(
string(
"Loaded program from ") + clbin +
".",
true);
603 device->store_cached_kernel(
program, cache_key, cache_locker);
606 add_log(
string(
"OpenCL program ") + program_name +
" not found on disk.",
true);
607 cache_locker.unlock();
614 needs_compiling =
false;
620 void OpenCLDevice::OpenCLProgram::compile()
624 string device_md5 = device->device_md5_hash(kernel_build_options);
628 ustring cache_key(program_name + device_md5);
629 program = device->load_cached_kernel(cache_key, cache_locker);
633 add_log(
string(
"OpenCL program ") + program_name +
" not found in cache.",
true);
636 string source = get_program_source(kernel_file);
638 string basename =
"cycles_kernel_" + program_name +
"_" + device_md5 +
"_" +
644 string clsrc, *debug_src =
NULL;
646 if (OpenCLInfo::use_debug()) {
651 if (
DebugFlags().running_inside_blender && compile_separate(clbin)) {
652 add_log(
string(
"Built and loaded program from ") + clbin +
".",
true);
657 add_log(
string(
"Separate-process building of ") + clbin +
658 " failed, will fall back to regular building.",
663 if (!compile_kernel(debug_src)) {
664 needs_compiling =
false;
669 if (!save_binary(clbin)) {
670 add_log(
string(
"Saving compiled OpenCL kernel to ") + clbin +
" failed!",
true);
675 device->store_cached_kernel(
program, cache_key, cache_locker);
679 needs_compiling =
false;
683 void OpenCLDevice::OpenCLProgram::create_kernels()
685 for (map<ustring, cl_kernel>::iterator kernel = kernels.begin(); kernel != kernels.end();
687 assert(kernel->second ==
NULL);
689 string name =
"kernel_ocl_" + kernel->first.string();
690 kernel->second = clCreateKernel(
program, name.c_str(), &ciErr);
691 if (device->opencl_error(ciErr)) {
692 add_error(
string(
"Error getting kernel ") + name +
" from program " + program_name +
": " +
693 clewErrorString(ciErr));
699 bool OpenCLDevice::OpenCLProgram::wait_for_availability()
701 add_log(
string(
"Waiting for availability of ") + program_name +
".",
true);
702 while (needs_compiling) {
708 void OpenCLDevice::OpenCLProgram::report_error()
717 cerr << error_msg << endl;
718 if (!compile_output.empty()) {
719 cerr <<
"OpenCL kernel build output for " << program_name <<
":" << endl;
720 cerr << compile_output << endl;
726 assert(kernels.size() == 1);
727 return kernels.begin()->second;
732 assert(kernels.count(name));
733 return kernels[name];
736 cl_device_type OpenCLInfo::device_type()
742 return CL_DEVICE_TYPE_ALL;
744 return CL_DEVICE_TYPE_DEFAULT;
746 return CL_DEVICE_TYPE_CPU;
748 return CL_DEVICE_TYPE_GPU;
750 return CL_DEVICE_TYPE_ACCELERATOR;
752 return CL_DEVICE_TYPE_ALL;
756 bool OpenCLInfo::use_debug()
761 bool OpenCLInfo::device_supported(
const string &platform_name,
const cl_device_id device_id)
763 cl_device_type device_type;
764 if (!get_device_type(device_id, &device_type)) {
768 if (!get_device_name(device_id, &device_name)) {
772 int driver_major = 0;
773 int driver_minor = 0;
774 if (!get_driver_version(device_id, &driver_major, &driver_minor)) {
777 VLOG(3) <<
"OpenCL driver version " << driver_major <<
"." << driver_minor;
779 if (getenv(
"CYCLES_OPENCL_TEST")) {
784 if (platform_name.find(
"Intel") != string::npos) {
785 if (device_type != CL_DEVICE_TYPE_GPU) {
796 if (device_name.find(
"Iris") != string::npos || device_name.find(
"Xe") != string::npos) {
802 if (platform_name ==
"AMD Accelerated Parallel Processing" &&
803 device_type == CL_DEVICE_TYPE_GPU) {
804 if (driver_major < 2236) {
805 VLOG(1) <<
"AMD driver version " << driver_major <<
"." << driver_minor <<
" not supported.";
808 const char *blacklist[] = {
815 for (
int i = 0; blacklist[i] !=
NULL; i++) {
816 if (device_name == blacklist[i]) {
817 VLOG(1) <<
"AMD device " << device_name <<
" not supported";
823 if (platform_name ==
"Apple" && device_type == CL_DEVICE_TYPE_GPU) {
829 bool OpenCLInfo::platform_version_check(cl_platform_id platform,
string *
error)
831 const int req_major = 1, req_minor = 1;
834 clGetPlatformInfo(platform, CL_PLATFORM_VERSION,
sizeof(version), &version,
NULL);
835 if (sscanf(version,
"OpenCL %d.%d", &major, &minor) < 2) {
837 *
error =
string_printf(
"OpenCL: failed to parse platform version string (%s).", version);
841 if (!((major == req_major && minor >= req_minor) || (major > req_major))) {
844 "OpenCL: platform version 1.1 or later required, found %d.%d", major, minor);
854 bool OpenCLInfo::get_device_version(cl_device_id device,
int *r_major,
int *r_minor,
string *
error)
857 clGetDeviceInfo(device, CL_DEVICE_OPENCL_C_VERSION,
sizeof(version), &version,
NULL);
858 if (sscanf(version,
"OpenCL C %d.%d", r_major, r_minor) < 2) {
860 *
error =
string_printf(
"OpenCL: failed to parse OpenCL C version string (%s).", version);
870 bool OpenCLInfo::device_version_check(cl_device_id device,
string *
error)
872 const int req_major = 1, req_minor = 1;
874 if (!get_device_version(device, &major, &minor,
error)) {
878 if (!((major == req_major && minor >= req_minor) || (major > req_major))) {
880 *
error =
string_printf(
"OpenCL: C version 1.1 or later required, found %d.%d", major, minor);
890 string OpenCLInfo::get_hardware_id(
const string &platform_name, cl_device_id device_id)
892 if (platform_name ==
"AMD Accelerated Parallel Processing" || platform_name ==
"Apple") {
894 cl_char topology[24];
895 if (clGetDeviceInfo(device_id, 0x4037,
sizeof(topology), topology,
NULL) == CL_SUCCESS &&
898 (
unsigned int)topology[21],
899 (
unsigned int)topology[22],
900 (
unsigned int)topology[23]);
903 else if (platform_name ==
"NVIDIA CUDA") {
905 cl_int bus_id, slot_id;
906 if (clGetDeviceInfo(device_id, 0x4008,
sizeof(cl_int), &bus_id,
NULL) == CL_SUCCESS &&
907 clGetDeviceInfo(device_id, 0x4009,
sizeof(cl_int), &slot_id,
NULL) == CL_SUCCESS) {
909 (
unsigned int)(bus_id),
910 (
unsigned int)(slot_id >> 3),
911 (
unsigned int)(slot_id & 0x7));
920 const cl_device_type device_type = OpenCLInfo::device_type();
921 static bool first_time =
true;
922 # define FIRST_VLOG(severity) \
926 usable_devices->clear();
928 if (device_type == 0) {
929 FIRST_VLOG(2) <<
"OpenCL devices are forced to be disabled.";
939 if (!get_platforms(&platform_ids, &
error)) {
940 FIRST_VLOG(2) <<
"Error fetching platforms:" << string(clewErrorString(
error));
944 if (platform_ids.size() == 0) {
945 FIRST_VLOG(2) <<
"No OpenCL platforms were found.";
950 for (
int platform = 0; platform < platform_ids.size(); platform++) {
951 cl_platform_id platform_id = platform_ids[platform];
952 string platform_name;
953 if (!get_platform_name(platform_id, &platform_name)) {
954 FIRST_VLOG(2) <<
"Failed to get platform name, ignoring.";
957 FIRST_VLOG(2) <<
"Enumerating devices for platform " << platform_name <<
".";
958 if (!platform_version_check(platform_id)) {
959 FIRST_VLOG(2) <<
"Ignoring platform " << platform_name
960 <<
" due to too old compiler version.";
963 if (!get_platform_devices(platform_id, device_type, &device_ids, &
error)) {
964 FIRST_VLOG(2) <<
"Ignoring platform " << platform_name
965 <<
", failed to fetch of devices: " << string(clewErrorString(
error));
968 if (device_ids.size() == 0) {
969 FIRST_VLOG(2) <<
"Ignoring platform " << platform_name <<
", it has no devices.";
972 for (
int num = 0; num < device_ids.size(); num++) {
973 const cl_device_id device_id = device_ids[num];
975 if (!get_device_name(device_id, &device_name, &
error)) {
976 FIRST_VLOG(2) <<
"Failed to fetch device name: " << string(clewErrorString(
error))
980 if (!device_version_check(device_id)) {
981 FIRST_VLOG(2) <<
"Ignoring device " << device_name <<
" due to old compiler version.";
984 if (device_supported(platform_name, device_id)) {
985 cl_device_type device_type;
986 if (!get_device_type(device_id, &device_type, &
error)) {
987 FIRST_VLOG(2) <<
"Ignoring device " << device_name
988 <<
", failed to fetch device type:" << string(clewErrorString(
error));
991 string readable_device_name = get_readable_device_name(device_id);
992 if (readable_device_name != device_name) {
993 FIRST_VLOG(2) <<
"Using more readable device name: " << readable_device_name;
995 FIRST_VLOG(2) <<
"Adding new device " << readable_device_name <<
".";
996 string hardware_id = get_hardware_id(platform_name, device_id);
997 string device_extensions = get_device_extensions(device_id);
998 usable_devices->push_back(OpenCLPlatformDevice(platform_id,
1002 readable_device_name,
1004 device_extensions));
1007 FIRST_VLOG(2) <<
"Ignoring device " << device_name <<
", not officially supported yet.";
1017 platform_ids->resize(0);
1018 cl_uint num_platforms;
1019 if (!get_num_platforms(&num_platforms,
error)) {
1024 platform_ids->resize(num_platforms);
1025 if ((
err = clGetPlatformIDs(num_platforms, &platform_ids->at(0),
NULL)) != CL_SUCCESS) {
1032 *
error = CL_SUCCESS;
1040 get_platforms(&platform_ids);
1041 return platform_ids;
1044 bool OpenCLInfo::get_num_platforms(cl_uint *num_platforms, cl_int *
error)
1047 if ((
err = clGetPlatformIDs(0,
NULL, num_platforms)) != CL_SUCCESS) {
1055 *
error = CL_SUCCESS;
1060 cl_uint OpenCLInfo::get_num_platforms()
1062 cl_uint num_platforms;
1063 if (!get_num_platforms(&num_platforms)) {
1066 return num_platforms;
1069 bool OpenCLInfo::get_platform_name(cl_platform_id platform_id,
string *platform_name)
1072 if (clGetPlatformInfo(platform_id, CL_PLATFORM_NAME,
sizeof(
buffer), &
buffer,
NULL) !=
1074 *platform_name =
"";
1081 string OpenCLInfo::get_platform_name(cl_platform_id platform_id)
1083 string platform_name;
1084 if (!get_platform_name(platform_id, &platform_name)) {
1087 return platform_name;
1090 bool OpenCLInfo::get_num_platform_devices(cl_platform_id platform_id,
1091 cl_device_type device_type,
1092 cl_uint *num_devices,
1096 if ((
err = clGetDeviceIDs(platform_id, device_type, 0,
NULL, num_devices)) != CL_SUCCESS) {
1104 *
error = CL_SUCCESS;
1109 cl_uint OpenCLInfo::get_num_platform_devices(cl_platform_id platform_id,
1110 cl_device_type device_type)
1112 cl_uint num_devices;
1113 if (!get_num_platform_devices(platform_id, device_type, &num_devices)) {
1119 bool OpenCLInfo::get_platform_devices(cl_platform_id platform_id,
1120 cl_device_type device_type,
1125 device_ids->resize(0);
1127 cl_uint num_devices;
1128 if (!get_num_platform_devices(platform_id, device_type, &num_devices,
error)) {
1132 device_ids->resize(num_devices);
1134 if ((
err = clGetDeviceIDs(platform_id, device_type, num_devices, &device_ids->at(0),
NULL)) !=
1142 *
error = CL_SUCCESS;
1148 cl_device_type device_type)
1151 get_platform_devices(platform_id, device_type, &
devices);
1155 bool OpenCLInfo::get_device_name(cl_device_id device_id,
string *device_name, cl_int *
error)
1168 *
error = CL_SUCCESS;
1174 string OpenCLInfo::get_device_name(cl_device_id device_id)
1177 if (!get_device_name(device_id, &device_name)) {
1183 bool OpenCLInfo::get_device_extensions(cl_device_id device_id,
1184 string *device_extensions,
1187 size_t extension_length = 0;
1190 if ((
err = clGetDeviceInfo(device_id, CL_DEVICE_EXTENSIONS, 0, 0, &extension_length)) !=
1195 *device_extensions =
"";
1199 if ((
err = clGetDeviceInfo(
1200 device_id, CL_DEVICE_EXTENSIONS, extension_length,
buffer.data(),
NULL)) !=
1205 *device_extensions =
"";
1209 *
error = CL_SUCCESS;
1211 *device_extensions = string(
buffer.data());
1215 string OpenCLInfo::get_device_extensions(cl_device_id device_id)
1217 string device_extensions;
1218 if (!get_device_extensions(device_id, &device_extensions)) {
1221 return device_extensions;
1224 bool OpenCLInfo::get_device_type(cl_device_id device_id,
1225 cl_device_type *device_type,
1229 if ((
err = clGetDeviceInfo(
1230 device_id, CL_DEVICE_TYPE,
sizeof(cl_device_type), device_type,
NULL)) != CL_SUCCESS) {
1238 *
error = CL_SUCCESS;
1243 cl_device_type OpenCLInfo::get_device_type(cl_device_id device_id)
1245 cl_device_type device_type;
1246 if (!get_device_type(device_id, &device_type)) {
1252 string OpenCLInfo::get_readable_device_name(cl_device_id device_id)
1255 char board_name[1024];
1257 if (clGetDeviceInfo(
1258 device_id, CL_DEVICE_BOARD_NAME_AMD,
sizeof(board_name), &board_name, &
length) ==
1260 if (
length != 0 && board_name[0] !=
'\0') {
1267 name = get_device_name(device_id);
1273 if (name ==
"Radeon RX Vega") {
1274 cl_int max_compute_units = 0;
1275 if (clGetDeviceInfo(device_id,
1276 CL_DEVICE_MAX_COMPUTE_UNITS,
1277 sizeof(max_compute_units),
1279 NULL) == CL_SUCCESS) {
1280 name +=
" " +
to_string(max_compute_units);
1285 if (get_device_type(device_id) & CL_DEVICE_TYPE_CPU) {
1286 name +=
" (OpenCL)";
1292 bool OpenCLInfo::get_driver_version(cl_device_id device_id,
int *major,
int *minor, cl_int *
error)
1296 if ((
err = clGetDeviceInfo(device_id, CL_DRIVER_VERSION,
sizeof(
buffer), &
buffer,
NULL)) !=
1304 *
error = CL_SUCCESS;
1306 if (sscanf(
buffer,
"%d.%d", major, minor) < 2) {
1313 int OpenCLInfo::mem_sub_ptr_alignment(cl_device_id device_id)
1315 int base_align_bits;
1316 if (clGetDeviceInfo(
1317 device_id, CL_DEVICE_MEM_BASE_ADDR_ALIGN,
sizeof(
int), &base_align_bits,
NULL) ==
1319 return base_align_bits / 8;
static DBVT_INLINE btScalar size(const btDbvtVolume &a)
SIMD_FORCE_INLINE btScalar length(const btQuaternion &q)
Return the length of a quaternion.
static char * basename(char *string)
bool device_opencl_compile_kernel(const vector< string > ¶meters)
#define CCL_NAMESPACE_END
__kernel void ccl_constant KernelData ccl_global void ccl_global char ccl_global int ccl_global char ccl_global unsigned int ccl_global float * buffer
double parameters[NUM_PARAMETERS]
static void error(const char *str)
int load(istream &in, Vec3r &v)
INLINE Rall1d< T, V, S > log(const Rall1d< T, V, S > &arg)
Vector< CPUDevice > devices
list of all CPUDevices. for every hardware thread an instance of CPUDevice is created
struct blender::compositor::@172::@175 opencl
std::string to_string(const T &n)
struct SELECTID_Context context
DebugFlags & DebugFlags()
string util_md5_string(const string &str)
string path_cache_get(const string &sub)
string path_get(const string &sub)
string path_source_replace_includes(const string &source, const string &path, const string &source_filename)
string path_files_md5_hash(const string &dir)
string path_join(const string &dir, const string &file)
bool path_exists(const string &path)
bool path_write_binary(const string &path, const vector< uint8_t > &binary)
bool path_write_text(const string &path, string &text)
bool path_read_binary(const string &path, vector< uint8_t > &binary)
CCL_NAMESPACE_BEGIN string string_printf(const char *format,...)
void string_replace(string &haystack, const string &needle, const string &other)
size_t system_physical_ram()
bool system_call_self(const vector< string > &args)
std::unique_lock< std::mutex > thread_scoped_lock
CCL_NAMESPACE_BEGIN typedef std::mutex thread_mutex
void time_sleep(double t)
CCL_NAMESPACE_BEGIN double time_dt()