Welcome to mirror list, hosted at ThFree Co, Russian Federation.

git.blender.org/blender.git - Unnamed repository; edit this file 'description' to name the repository.
summaryrefslogtreecommitdiff
path: root/intern/cycles/device
diff options
context:
space:
mode:
Diffstat (limited to 'intern/cycles/device')
-rw-r--r--intern/cycles/device/device.cpp2
-rw-r--r--intern/cycles/device/device.h10
-rw-r--r--intern/cycles/device/device_split_kernel.cpp4
-rw-r--r--intern/cycles/device/device_split_kernel.h1
-rw-r--r--intern/cycles/device/opencl/opencl.h5
-rw-r--r--intern/cycles/device/opencl/opencl_split.cpp63
-rw-r--r--intern/cycles/device/opencl/opencl_util.cpp51
7 files changed, 131 insertions, 5 deletions
diff --git a/intern/cycles/device/device.cpp b/intern/cycles/device/device.cpp
index 949c5f932a4..31671e76ec3 100644
--- a/intern/cycles/device/device.cpp
+++ b/intern/cycles/device/device.cpp
@@ -69,6 +69,8 @@ std::ostream& operator <<(std::ostream &os,
<< string_from_bool(requested_features.use_transparent) << std::endl;
os << "Use Principled BSDF: "
<< string_from_bool(requested_features.use_principled) << std::endl;
+ os << "Use Denoising: "
+ << string_from_bool(requested_features.use_denoising) << std::endl;
return os;
}
diff --git a/intern/cycles/device/device.h b/intern/cycles/device/device.h
index c22969d7dc6..68a555c1a93 100644
--- a/intern/cycles/device/device.h
+++ b/intern/cycles/device/device.h
@@ -127,6 +127,9 @@ public:
/* Per-uber shader usage flags. */
bool use_principled;
+ /* Denoising features. */
+ bool use_denoising;
+
DeviceRequestedFeatures()
{
/* TODO(sergey): Find more meaningful defaults. */
@@ -145,6 +148,7 @@ public:
use_transparent = false;
use_shadow_tricks = false;
use_principled = false;
+ use_denoising = false;
}
bool modified(const DeviceRequestedFeatures& requested_features)
@@ -163,7 +167,8 @@ public:
use_patch_evaluation == requested_features.use_patch_evaluation &&
use_transparent == requested_features.use_transparent &&
use_shadow_tricks == requested_features.use_shadow_tricks &&
- use_principled == requested_features.use_principled);
+ use_principled == requested_features.use_principled &&
+ use_denoising == requested_features.use_denoising);
}
/* Convert the requested features structure to a build options,
@@ -213,6 +218,9 @@ public:
if(!use_principled) {
build_options += " -D__NO_PRINCIPLED__";
}
+ if(!use_denoising) {
+ build_options += " -D__NO_DENOISING__";
+ }
return build_options;
}
};
diff --git a/intern/cycles/device/device_split_kernel.cpp b/intern/cycles/device/device_split_kernel.cpp
index dddd19f179f..d2b3a89fa98 100644
--- a/intern/cycles/device/device_split_kernel.cpp
+++ b/intern/cycles/device/device_split_kernel.cpp
@@ -47,6 +47,7 @@ DeviceSplitKernel::DeviceSplitKernel(Device *device) : device(device)
kernel_direct_lighting = NULL;
kernel_shadow_blocked_ao = NULL;
kernel_shadow_blocked_dl = NULL;
+ kernel_enqueue_inactive = NULL;
kernel_next_iteration_setup = NULL;
kernel_indirect_subsurface = NULL;
kernel_buffer_update = NULL;
@@ -74,6 +75,7 @@ DeviceSplitKernel::~DeviceSplitKernel()
delete kernel_direct_lighting;
delete kernel_shadow_blocked_ao;
delete kernel_shadow_blocked_dl;
+ delete kernel_enqueue_inactive;
delete kernel_next_iteration_setup;
delete kernel_indirect_subsurface;
delete kernel_buffer_update;
@@ -101,6 +103,7 @@ bool DeviceSplitKernel::load_kernels(const DeviceRequestedFeatures& requested_fe
LOAD_KERNEL(direct_lighting);
LOAD_KERNEL(shadow_blocked_ao);
LOAD_KERNEL(shadow_blocked_dl);
+ LOAD_KERNEL(enqueue_inactive);
LOAD_KERNEL(next_iteration_setup);
LOAD_KERNEL(indirect_subsurface);
LOAD_KERNEL(buffer_update);
@@ -256,6 +259,7 @@ bool DeviceSplitKernel::path_trace(DeviceTask *task,
ENQUEUE_SPLIT_KERNEL(direct_lighting, global_size, local_size);
ENQUEUE_SPLIT_KERNEL(shadow_blocked_ao, global_size, local_size);
ENQUEUE_SPLIT_KERNEL(shadow_blocked_dl, global_size, local_size);
+ ENQUEUE_SPLIT_KERNEL(enqueue_inactive, global_size, local_size);
ENQUEUE_SPLIT_KERNEL(next_iteration_setup, global_size, local_size);
ENQUEUE_SPLIT_KERNEL(indirect_subsurface, global_size, local_size);
ENQUEUE_SPLIT_KERNEL(queue_enqueue, global_size, local_size);
diff --git a/intern/cycles/device/device_split_kernel.h b/intern/cycles/device/device_split_kernel.h
index 68c2ba974a5..2bac1998cb7 100644
--- a/intern/cycles/device/device_split_kernel.h
+++ b/intern/cycles/device/device_split_kernel.h
@@ -69,6 +69,7 @@ private:
SplitKernelFunction *kernel_direct_lighting;
SplitKernelFunction *kernel_shadow_blocked_ao;
SplitKernelFunction *kernel_shadow_blocked_dl;
+ SplitKernelFunction *kernel_enqueue_inactive;
SplitKernelFunction *kernel_next_iteration_setup;
SplitKernelFunction *kernel_indirect_subsurface;
SplitKernelFunction *kernel_buffer_update;
diff --git a/intern/cycles/device/opencl/opencl.h b/intern/cycles/device/opencl/opencl.h
index 52851061d7b..399fae9b42e 100644
--- a/intern/cycles/device/opencl/opencl.h
+++ b/intern/cycles/device/opencl/opencl.h
@@ -130,6 +130,11 @@ public:
cl_int* error = NULL);
static cl_device_type get_device_type(cl_device_id device_id);
+ static bool get_driver_version(cl_device_id device_id,
+ int *major,
+ int *minor,
+ cl_int* error = NULL);
+
static int mem_address_alignment(cl_device_id device_id);
/* Get somewhat more readable device name.
diff --git a/intern/cycles/device/opencl/opencl_split.cpp b/intern/cycles/device/opencl/opencl_split.cpp
index 76dcbd6fc9a..08b632ee9d3 100644
--- a/intern/cycles/device/opencl/opencl_split.cpp
+++ b/intern/cycles/device/opencl/opencl_split.cpp
@@ -176,17 +176,62 @@ protected:
friend class OpenCLSplitKernelFunction;
};
+struct CachedSplitMemory {
+ int id;
+ device_memory *split_data;
+ device_memory *ray_state;
+ device_ptr *rng_state;
+ device_memory *queue_index;
+ device_memory *use_queues_flag;
+ device_memory *work_pools;
+ device_ptr *buffer;
+};
+
class OpenCLSplitKernelFunction : public SplitKernelFunction {
public:
OpenCLDeviceSplitKernel* device;
OpenCLDeviceBase::OpenCLProgram program;
+ CachedSplitMemory& cached_memory;
+ int cached_id;
+
+ OpenCLSplitKernelFunction(OpenCLDeviceSplitKernel* device, CachedSplitMemory& cached_memory) :
+ device(device), cached_memory(cached_memory), cached_id(cached_memory.id-1)
+ {
+ }
- OpenCLSplitKernelFunction(OpenCLDeviceSplitKernel* device) : device(device) {}
- ~OpenCLSplitKernelFunction() { program.release(); }
+ ~OpenCLSplitKernelFunction()
+ {
+ program.release();
+ }
virtual bool enqueue(const KernelDimensions& dim, device_memory& kg, device_memory& data)
{
- device->kernel_set_args(program(), 0, kg, data);
+ if(cached_id != cached_memory.id) {
+ cl_uint start_arg_index =
+ device->kernel_set_args(program(),
+ 0,
+ kg,
+ data,
+ *cached_memory.split_data,
+ *cached_memory.ray_state,
+ *cached_memory.rng_state);
+
+/* TODO(sergey): Avoid map lookup here. */
+#define KERNEL_TEX(type, ttype, name) \
+ device->set_kernel_arg_mem(program(), &start_arg_index, #name);
+#include "kernel/kernel_textures.h"
+#undef KERNEL_TEX
+
+ start_arg_index +=
+ device->kernel_set_args(program(),
+ start_arg_index,
+ *cached_memory.queue_index,
+ *cached_memory.use_queues_flag,
+ *cached_memory.work_pools,
+ *cached_memory.buffer);
+
+ cached_id = cached_memory.id;
+ }
device->ciErr = clEnqueueNDRangeKernel(device->cqCommandQueue,
program(),
@@ -213,6 +258,7 @@ public:
class OpenCLSplitKernel : public DeviceSplitKernel {
OpenCLDeviceSplitKernel *device;
+ CachedSplitMemory cached_memory;
public:
explicit OpenCLSplitKernel(OpenCLDeviceSplitKernel *device) : DeviceSplitKernel(device), device(device) {
}
@@ -220,7 +266,7 @@ public:
virtual SplitKernelFunction* get_split_kernel_function(string kernel_name,
const DeviceRequestedFeatures& requested_features)
{
- OpenCLSplitKernelFunction* kernel = new OpenCLSplitKernelFunction(device);
+ OpenCLSplitKernelFunction* kernel = new OpenCLSplitKernelFunction(device, cached_memory);
bool single_program = OpenCLInfo::use_single_program();
kernel->program =
@@ -349,6 +395,15 @@ public:
return false;
}
+ cached_memory.split_data = &split_data;
+ cached_memory.ray_state = &ray_state;
+ cached_memory.rng_state = &rtile.rng_state;
+ cached_memory.queue_index = &queue_index;
+ cached_memory.use_queues_flag = &use_queues_flag;
+ cached_memory.work_pools = &work_pool_wgs;
+ cached_memory.buffer = &rtile.buffer;
+ cached_memory.id++;
+
return true;
}
diff --git a/intern/cycles/device/opencl/opencl_util.cpp b/intern/cycles/device/opencl/opencl_util.cpp
index d27aa05c312..8ba2a8e26da 100644
--- a/intern/cycles/device/opencl/opencl_util.cpp
+++ b/intern/cycles/device/opencl/opencl_util.cpp
@@ -608,6 +608,14 @@ bool OpenCLInfo::device_supported(const string& platform_name,
if(!get_device_name(device_id, &device_name)) {
return false;
}
+
+ int driver_major = 0;
+ int driver_minor = 0;
+ if(!get_driver_version(device_id, &driver_major, &driver_minor)) {
+ return false;
+ }
+ VLOG(3) << "OpenCL driver version " << driver_major << "." << driver_minor;
+
/* It is possible tyo have Iris GPU on AMD/Apple OpenCL framework
* (aka, it will not be on Intel framework). This isn't supported
* and needs an explicit blacklist.
@@ -618,6 +626,21 @@ bool OpenCLInfo::device_supported(const string& platform_name,
if(platform_name == "AMD Accelerated Parallel Processing" &&
device_type == CL_DEVICE_TYPE_GPU)
{
+ if(driver_major < 2236) {
+ VLOG(1) << "AMD driver version " << driver_major << "." << driver_minor << " not supported.";
+ return false;
+ }
+ const char *blacklist[] = {
+ /* GCN 1 */
+ "Tahiti", "Pitcairn", "Capeverde", "Oland",
+ NULL
+ };
+ for (int i = 0; blacklist[i] != NULL; i++) {
+ if(device_name == blacklist[i]) {
+ VLOG(1) << "AMD device " << device_name << " not supported";
+ return false;
+ }
+ }
return true;
}
if(platform_name == "Apple" && device_type == CL_DEVICE_TYPE_GPU) {
@@ -1073,6 +1096,34 @@ string OpenCLInfo::get_readable_device_name(cl_device_id device_id)
return get_device_name(device_id);
}
+bool OpenCLInfo::get_driver_version(cl_device_id device_id,
+ int *major,
+ int *minor,
+ cl_int* error)
+{
+ char buffer[1024];
+ cl_int err;
+ if((err = clGetDeviceInfo(device_id,
+ CL_DRIVER_VERSION,
+ sizeof(buffer),
+ &buffer,
+ NULL)) != CL_SUCCESS)
+ {
+ if(error != NULL) {
+ *error = err;
+ }
+ return false;
+ }
+ if(error != NULL) {
+ *error = CL_SUCCESS;
+ }
+ if(sscanf(buffer, "%d.%d", major, minor) < 2) {
+ VLOG(1) << string_printf("OpenCL: failed to parse driver version string (%s).", buffer);
+ return false;
+ }
+ return true;
+}
+
int OpenCLInfo::mem_address_alignment(cl_device_id device_id)
{
int base_align_bits;
generated by cgit v1.2.3 (git 2.25.1) at 2024-07-12 21:30:27 +0300