Cycles OpenCL: Remove OpenCL MegaKernel
Using OpenCL MegaKernel has been slow and therefore not usefull. This patch will remove the mega kernel from the OpenCL codebase and the OpenCLDeviceBase class. T61736: removal of mega kernel T61703: baking does not work with mega kernel Tags: #cycles Differential Revision: https://developer.blender.org/D4383
This commit is contained in:
parent
667033e89e
commit
949ab753bb
Notes:
blender-bot
2023-02-14 00:06:52 +01:00
Referenced by commit 44c6b6615b
, OpenCL: Bring back CYCLES_OPENCL_TEST override
Referenced by issue #73734, Can't force the detection of OpenCL devices
|
@ -1634,7 +1634,6 @@ class CYCLES_RENDER_PT_debug(CyclesButtonsPanel, Panel):
|
|||
|
||||
col = layout.column()
|
||||
col.label('OpenCL Flags:')
|
||||
col.prop(cscene, "debug_opencl_kernel_type", text="Kernel")
|
||||
col.prop(cscene, "debug_opencl_device_type", text="Device")
|
||||
col.prop(cscene, "debug_opencl_kernel_single_program", text="Single Program")
|
||||
col.prop(cscene, "debug_use_opencl_debug", text="Debug")
|
||||
|
|
|
@ -67,7 +67,6 @@ bool debug_flags_sync_from_scene(BL::Scene b_scene)
|
|||
PointerRNA cscene = RNA_pointer_get(&b_scene.ptr, "cycles");
|
||||
/* Backup some settings for comparison. */
|
||||
DebugFlags::OpenCL::DeviceType opencl_device_type = flags.opencl.device_type;
|
||||
DebugFlags::OpenCL::KernelType opencl_kernel_type = flags.opencl.kernel_type;
|
||||
/* Synchronize shared flags. */
|
||||
flags.viewport_static_bvh = get_enum(cscene, "debug_bvh_type");
|
||||
/* Synchronize CPU flags. */
|
||||
|
@ -81,18 +80,6 @@ bool debug_flags_sync_from_scene(BL::Scene b_scene)
|
|||
/* Synchronize CUDA flags. */
|
||||
flags.cuda.adaptive_compile = get_boolean(cscene, "debug_use_cuda_adaptive_compile");
|
||||
flags.cuda.split_kernel = get_boolean(cscene, "debug_use_cuda_split_kernel");
|
||||
/* Synchronize OpenCL kernel type. */
|
||||
switch(get_enum(cscene, "debug_opencl_kernel_type")) {
|
||||
case 0:
|
||||
flags.opencl.kernel_type = DebugFlags::OpenCL::KERNEL_DEFAULT;
|
||||
break;
|
||||
case 1:
|
||||
flags.opencl.kernel_type = DebugFlags::OpenCL::KERNEL_MEGA;
|
||||
break;
|
||||
case 2:
|
||||
flags.opencl.kernel_type = DebugFlags::OpenCL::KERNEL_SPLIT;
|
||||
break;
|
||||
}
|
||||
/* Synchronize OpenCL device type. */
|
||||
switch(get_enum(cscene, "debug_opencl_device_type")) {
|
||||
case 0:
|
||||
|
@ -118,8 +105,7 @@ bool debug_flags_sync_from_scene(BL::Scene b_scene)
|
|||
flags.opencl.debug = get_boolean(cscene, "debug_use_opencl_debug");
|
||||
flags.opencl.mem_limit = ((size_t)get_int(cscene, "debug_opencl_mem_limit"))*1024*1024;
|
||||
flags.opencl.single_program = get_boolean(cscene, "debug_opencl_kernel_single_program");
|
||||
return flags.opencl.device_type != opencl_device_type ||
|
||||
flags.opencl.kernel_type != opencl_kernel_type;
|
||||
return flags.opencl.device_type != opencl_device_type;
|
||||
}
|
||||
|
||||
/* Reset debug flags to default values.
|
||||
|
@ -130,10 +116,8 @@ bool debug_flags_reset()
|
|||
DebugFlagsRef flags = DebugFlags();
|
||||
/* Backup some settings for comparison. */
|
||||
DebugFlags::OpenCL::DeviceType opencl_device_type = flags.opencl.device_type;
|
||||
DebugFlags::OpenCL::KernelType opencl_kernel_type = flags.opencl.kernel_type;
|
||||
flags.reset();
|
||||
return flags.opencl.device_type != opencl_device_type ||
|
||||
flags.opencl.kernel_type != opencl_kernel_type;
|
||||
return flags.opencl.device_type != opencl_device_type;
|
||||
}
|
||||
|
||||
} /* namespace */
|
||||
|
|
|
@ -37,8 +37,6 @@ set(SRC_OPENCL
|
|||
opencl/opencl.h
|
||||
opencl/memory_manager.h
|
||||
|
||||
opencl/opencl_base.cpp
|
||||
opencl/opencl_mega.cpp
|
||||
opencl/opencl_split.cpp
|
||||
opencl/opencl_util.cpp
|
||||
opencl/memory_manager.cpp
|
||||
|
|
|
@ -29,19 +29,7 @@ CCL_NAMESPACE_BEGIN
|
|||
|
||||
Device *device_opencl_create(DeviceInfo& info, Stats &stats, Profiler &profiler, bool background)
|
||||
{
|
||||
vector<OpenCLPlatformDevice> usable_devices;
|
||||
OpenCLInfo::get_usable_devices(&usable_devices);
|
||||
assert(info.num < usable_devices.size());
|
||||
const OpenCLPlatformDevice& platform_device = usable_devices[info.num];
|
||||
const string& platform_name = platform_device.platform_name;
|
||||
const cl_device_type device_type = platform_device.device_type;
|
||||
if(OpenCLInfo::kernel_use_split(platform_name, device_type)) {
|
||||
VLOG(1) << "Using split kernel.";
|
||||
return opencl_create_split_device(info, stats, profiler, background);
|
||||
} else {
|
||||
VLOG(1) << "Using mega kernel.";
|
||||
return opencl_create_mega_device(info, stats, profiler, background);
|
||||
}
|
||||
return opencl_create_split_device(info, stats, profiler, background);
|
||||
}
|
||||
|
||||
bool device_opencl_init()
|
||||
|
@ -111,7 +99,6 @@ void device_opencl_info(vector<DeviceInfo>& devices)
|
|||
foreach(OpenCLPlatformDevice& platform_device, usable_devices) {
|
||||
/* Compute unique ID for persistent user preferences. */
|
||||
const string& platform_name = platform_device.platform_name;
|
||||
const cl_device_type device_type = platform_device.device_type;
|
||||
const string& device_name = platform_device.device_name;
|
||||
string hardware_id = platform_device.hardware_id;
|
||||
if(hardware_id == "") {
|
||||
|
@ -133,8 +120,7 @@ void device_opencl_info(vector<DeviceInfo>& devices)
|
|||
/* We don't know if it's used for display, but assume it is. */
|
||||
info.display_device = true;
|
||||
info.advanced_shading = OpenCLInfo::kernel_use_advanced_shading(platform_name);
|
||||
info.use_split_kernel = OpenCLInfo::kernel_use_split(platform_name,
|
||||
device_type);
|
||||
info.use_split_kernel = true;
|
||||
info.has_volume_decoupled = false;
|
||||
info.id = id;
|
||||
|
||||
|
|
|
@ -28,7 +28,7 @@ void MemoryManager::DeviceBuffer::add_allocation(Allocation& allocation)
|
|||
allocations.push_back(&allocation);
|
||||
}
|
||||
|
||||
void MemoryManager::DeviceBuffer::update_device_memory(OpenCLDeviceBase *device)
|
||||
void MemoryManager::DeviceBuffer::update_device_memory(OpenCLDevice *device)
|
||||
{
|
||||
bool need_realloc = false;
|
||||
|
||||
|
@ -142,7 +142,7 @@ void MemoryManager::DeviceBuffer::update_device_memory(OpenCLDeviceBase *device)
|
|||
clFinish(device->cqCommandQueue);
|
||||
}
|
||||
|
||||
void MemoryManager::DeviceBuffer::free(OpenCLDeviceBase *)
|
||||
void MemoryManager::DeviceBuffer::free(OpenCLDevice *)
|
||||
{
|
||||
buffer->free();
|
||||
}
|
||||
|
@ -160,7 +160,7 @@ MemoryManager::DeviceBuffer* MemoryManager::smallest_device_buffer()
|
|||
return smallest;
|
||||
}
|
||||
|
||||
MemoryManager::MemoryManager(OpenCLDeviceBase *device)
|
||||
MemoryManager::MemoryManager(OpenCLDevice *device)
|
||||
: device(device), need_update(false)
|
||||
{
|
||||
foreach(DeviceBuffer& device_buffer, device_buffers) {
|
||||
|
|
|
@ -26,7 +26,7 @@
|
|||
|
||||
CCL_NAMESPACE_BEGIN
|
||||
|
||||
class OpenCLDeviceBase;
|
||||
class OpenCLDevice;
|
||||
|
||||
class MemoryManager {
|
||||
public:
|
||||
|
@ -73,12 +73,12 @@ private:
|
|||
|
||||
void add_allocation(Allocation& allocation);
|
||||
|
||||
void update_device_memory(OpenCLDeviceBase *device);
|
||||
void update_device_memory(OpenCLDevice *device);
|
||||
|
||||
void free(OpenCLDeviceBase *device);
|
||||
void free(OpenCLDevice *device);
|
||||
};
|
||||
|
||||
OpenCLDeviceBase *device;
|
||||
OpenCLDevice *device;
|
||||
|
||||
DeviceBuffer device_buffers[NUM_DEVICE_BUFFERS];
|
||||
|
||||
|
@ -90,7 +90,7 @@ private:
|
|||
DeviceBuffer* smallest_device_buffer();
|
||||
|
||||
public:
|
||||
MemoryManager(OpenCLDeviceBase *device);
|
||||
MemoryManager(OpenCLDevice *device);
|
||||
|
||||
void free(); /* Free all memory. */
|
||||
|
||||
|
|
|
@ -18,6 +18,7 @@
|
|||
|
||||
#include "device/device.h"
|
||||
#include "device/device_denoising.h"
|
||||
#include "device/device_split_kernel.h"
|
||||
|
||||
#include "util/util_map.h"
|
||||
#include "util/util_param.h"
|
||||
|
@ -84,8 +85,6 @@ public:
|
|||
static cl_device_type device_type();
|
||||
static bool use_debug();
|
||||
static bool kernel_use_advanced_shading(const string& platform_name);
|
||||
static bool kernel_use_split(const string& platform_name,
|
||||
const cl_device_type device_type);
|
||||
static bool device_supported(const string& platform_name,
|
||||
const cl_device_id device_id);
|
||||
static bool platform_version_check(cl_platform_id platform,
|
||||
|
@ -259,7 +258,7 @@ public:
|
|||
} \
|
||||
} (void) 0
|
||||
|
||||
class OpenCLDeviceBase : public Device
|
||||
class OpenCLDevice : public Device
|
||||
{
|
||||
public:
|
||||
DedicatedTaskPool task_pool;
|
||||
|
@ -273,7 +272,7 @@ public:
|
|||
class OpenCLProgram {
|
||||
public:
|
||||
OpenCLProgram() : loaded(false), program(NULL), device(NULL) {}
|
||||
OpenCLProgram(OpenCLDeviceBase *device,
|
||||
OpenCLProgram(OpenCLDevice *device,
|
||||
const string& program_name,
|
||||
const string& kernel_name,
|
||||
const string& kernel_build_options,
|
||||
|
@ -311,7 +310,7 @@ public:
|
|||
|
||||
bool loaded;
|
||||
cl_program program;
|
||||
OpenCLDeviceBase *device;
|
||||
OpenCLDevice *device;
|
||||
|
||||
/* Used for the OpenCLCache key. */
|
||||
string program_name;
|
||||
|
@ -325,6 +324,32 @@ public:
|
|||
map<ustring, cl_kernel> kernels;
|
||||
};
|
||||
|
||||
DeviceSplitKernel *split_kernel;
|
||||
|
||||
OpenCLProgram program_data_init;
|
||||
OpenCLProgram program_state_buffer_size;
|
||||
|
||||
OpenCLProgram program_split;
|
||||
|
||||
OpenCLProgram program_path_init;
|
||||
OpenCLProgram program_scene_intersect;
|
||||
OpenCLProgram program_lamp_emission;
|
||||
OpenCLProgram program_do_volume;
|
||||
OpenCLProgram program_queue_enqueue;
|
||||
OpenCLProgram program_indirect_background;
|
||||
OpenCLProgram program_shader_setup;
|
||||
OpenCLProgram program_shader_sort;
|
||||
OpenCLProgram program_shader_eval;
|
||||
OpenCLProgram program_holdout_emission_blurring_pathtermination_ao;
|
||||
OpenCLProgram program_subsurface_scatter;
|
||||
OpenCLProgram program_direct_lighting;
|
||||
OpenCLProgram program_shadow_blocked_ao;
|
||||
OpenCLProgram program_shadow_blocked_dl;
|
||||
OpenCLProgram program_enqueue_inactive;
|
||||
OpenCLProgram program_next_iteration_setup;
|
||||
OpenCLProgram program_indirect_subsurface;
|
||||
OpenCLProgram program_buffer_update;
|
||||
|
||||
OpenCLProgram base_program;
|
||||
OpenCLProgram bake_program;
|
||||
OpenCLProgram displace_program;
|
||||
|
@ -346,8 +371,8 @@ public:
|
|||
void opencl_error(const string& message);
|
||||
void opencl_assert_err(cl_int err, const char* where);
|
||||
|
||||
OpenCLDeviceBase(DeviceInfo& info, Stats &stats, Profiler &profiler, bool background_);
|
||||
~OpenCLDeviceBase();
|
||||
OpenCLDevice(DeviceInfo& info, Stats &stats, Profiler &profiler, bool background_);
|
||||
~OpenCLDevice();
|
||||
|
||||
static void CL_CALLBACK context_notify_callback(const char *err_info,
|
||||
const void * /*private_info*/, size_t /*cb*/, void *user_data);
|
||||
|
@ -355,17 +380,14 @@ public:
|
|||
bool opencl_version_check();
|
||||
|
||||
string device_md5_hash(string kernel_custom_build_options = "");
|
||||
virtual bool load_kernels(const DeviceRequestedFeatures& requested_features);
|
||||
|
||||
/* Has to be implemented by the real device classes.
|
||||
* The base device will then load all these programs. */
|
||||
virtual bool add_kernel_programs(const DeviceRequestedFeatures& requested_features,
|
||||
vector<OpenCLProgram*> &programs) = 0;
|
||||
bool load_kernels(const DeviceRequestedFeatures& requested_features);
|
||||
|
||||
/* Get the name of the opencl program for the given kernel */
|
||||
virtual const string get_opencl_program_name(bool single_program, const string& kernel_name) = 0;
|
||||
const string get_opencl_program_name(bool single_program, const string& kernel_name);
|
||||
/* Get the program file name to compile (*.cl) for the given kernel */
|
||||
virtual const string get_opencl_program_filename(bool single_program, const string& kernel_name) = 0;
|
||||
const string get_opencl_program_filename(bool single_program, const string& kernel_name);
|
||||
string get_build_options(const DeviceRequestedFeatures& requested_features);
|
||||
string get_build_options_for_bake(const DeviceRequestedFeatures& requested_features);
|
||||
|
||||
void mem_alloc(device_memory& mem);
|
||||
void mem_copy_to(device_memory& mem);
|
||||
|
@ -393,10 +415,10 @@ public:
|
|||
|
||||
class OpenCLDeviceTask : public DeviceTask {
|
||||
public:
|
||||
OpenCLDeviceTask(OpenCLDeviceBase *device, DeviceTask& task)
|
||||
OpenCLDeviceTask(OpenCLDevice *device, DeviceTask& task)
|
||||
: DeviceTask(task)
|
||||
{
|
||||
run = function_bind(&OpenCLDeviceBase::thread_run,
|
||||
run = function_bind(&OpenCLDevice::thread_run,
|
||||
device,
|
||||
this);
|
||||
}
|
||||
|
@ -422,9 +444,16 @@ public:
|
|||
task_pool.cancel();
|
||||
}
|
||||
|
||||
virtual void thread_run(DeviceTask * /*task*/) = 0;
|
||||
void thread_run(DeviceTask *task);
|
||||
|
||||
virtual BVHLayoutMask get_bvh_layout_mask() const {
|
||||
return BVH_LAYOUT_BVH2;
|
||||
}
|
||||
|
||||
virtual bool show_samples() const {
|
||||
return true;
|
||||
}
|
||||
|
||||
virtual bool is_split_kernel() = 0;
|
||||
|
||||
protected:
|
||||
string kernel_build_options(const string *debug_src = NULL);
|
||||
|
@ -566,18 +595,15 @@ protected:
|
|||
|
||||
/* ** Those guys are for workign around some compiler-specific bugs ** */
|
||||
|
||||
virtual cl_program load_cached_kernel(
|
||||
cl_program load_cached_kernel(
|
||||
ustring key,
|
||||
thread_scoped_lock& cache_locker);
|
||||
|
||||
virtual void store_cached_kernel(
|
||||
void store_cached_kernel(
|
||||
cl_program program,
|
||||
ustring key,
|
||||
thread_scoped_lock& cache_locker);
|
||||
|
||||
virtual string build_options_for_bake_program(
|
||||
const DeviceRequestedFeatures& /*requested_features*/);
|
||||
|
||||
private:
|
||||
MemoryManager memory_manager;
|
||||
friend class MemoryManager;
|
||||
|
@ -592,9 +618,11 @@ private:
|
|||
|
||||
protected:
|
||||
void flush_texture_buffers();
|
||||
|
||||
friend class OpenCLSplitKernel;
|
||||
friend class OpenCLSplitKernelFunction;
|
||||
};
|
||||
|
||||
Device *opencl_create_mega_device(DeviceInfo& info, Stats& stats, Profiler &profiler, bool background);
|
||||
Device *opencl_create_split_device(DeviceInfo& info, Stats& stats, Profiler &profiler, bool background);
|
||||
|
||||
CCL_NAMESPACE_END
|
||||
|
|
File diff suppressed because it is too large
Load Diff
|
@ -1,186 +0,0 @@
|
|||
/*
|
||||
* Copyright 2011-2013 Blender Foundation
|
||||
*
|
||||
* Licensed under the Apache License, Version 2.0 (the "License");
|
||||
* you may not use this file except in compliance with the License.
|
||||
* You may obtain a copy of the License at
|
||||
*
|
||||
* http://www.apache.org/licenses/LICENSE-2.0
|
||||
*
|
||||
* Unless required by applicable law or agreed to in writing, software
|
||||
* distributed under the License is distributed on an "AS IS" BASIS,
|
||||
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
||||
* See the License for the specific language governing permissions and
|
||||
* limitations under the License.
|
||||
*/
|
||||
|
||||
#ifdef WITH_OPENCL
|
||||
|
||||
#include "device/opencl/opencl.h"
|
||||
|
||||
#include "render/buffers.h"
|
||||
|
||||
#include "kernel/kernel_types.h"
|
||||
|
||||
#include "util/util_md5.h"
|
||||
#include "util/util_path.h"
|
||||
#include "util/util_time.h"
|
||||
|
||||
CCL_NAMESPACE_BEGIN
|
||||
|
||||
class OpenCLDeviceMegaKernel : public OpenCLDeviceBase
|
||||
{
|
||||
public:
|
||||
OpenCLProgram path_trace_program;
|
||||
|
||||
OpenCLDeviceMegaKernel(DeviceInfo& info, Stats &stats, Profiler &profiler, bool background_)
|
||||
: OpenCLDeviceBase(info, stats, profiler, background_),
|
||||
path_trace_program(this,
|
||||
get_opencl_program_name(false, "megakernel"),
|
||||
get_opencl_program_filename(false, "megakernel"),
|
||||
"-D__COMPILE_ONLY_MEGAKERNEL__ ")
|
||||
{
|
||||
}
|
||||
|
||||
|
||||
virtual bool show_samples() const
|
||||
{
|
||||
return true;
|
||||
}
|
||||
|
||||
virtual BVHLayoutMask get_bvh_layout_mask() const
|
||||
{
|
||||
return BVH_LAYOUT_BVH2;
|
||||
}
|
||||
|
||||
const string get_opencl_program_name(bool /*single_program*/, const string& kernel_name)
|
||||
{
|
||||
return kernel_name;
|
||||
}
|
||||
|
||||
const string get_opencl_program_filename(bool /*single_program*/, const string& /*kernel_name*/)
|
||||
{
|
||||
return "kernel.cl";
|
||||
}
|
||||
|
||||
virtual bool add_kernel_programs(const DeviceRequestedFeatures& /*requested_features*/,
|
||||
vector<OpenCLProgram*> &programs)
|
||||
{
|
||||
path_trace_program.add_kernel(ustring("path_trace"));
|
||||
programs.push_back(&path_trace_program);
|
||||
return true;
|
||||
}
|
||||
|
||||
~OpenCLDeviceMegaKernel()
|
||||
{
|
||||
task_pool.stop();
|
||||
path_trace_program.release();
|
||||
}
|
||||
|
||||
void path_trace(RenderTile& rtile, int sample)
|
||||
{
|
||||
scoped_timer timer(&rtile.buffers->render_time);
|
||||
|
||||
/* 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(rtile.buffer);
|
||||
cl_int d_x = rtile.x;
|
||||
cl_int d_y = rtile.y;
|
||||
cl_int d_w = rtile.w;
|
||||
cl_int d_h = rtile.h;
|
||||
cl_int d_offset = rtile.offset;
|
||||
cl_int d_stride = rtile.stride;
|
||||
|
||||
/* Sample arguments. */
|
||||
cl_int d_sample = sample;
|
||||
|
||||
cl_kernel ckPathTraceKernel = path_trace_program(ustring("path_trace"));
|
||||
|
||||
cl_uint start_arg_index =
|
||||
kernel_set_args(ckPathTraceKernel,
|
||||
0,
|
||||
d_data,
|
||||
d_buffer);
|
||||
|
||||
set_kernel_arg_buffers(ckPathTraceKernel, &start_arg_index);
|
||||
|
||||
start_arg_index += kernel_set_args(ckPathTraceKernel,
|
||||
start_arg_index,
|
||||
d_sample,
|
||||
d_x,
|
||||
d_y,
|
||||
d_w,
|
||||
d_h,
|
||||
d_offset,
|
||||
d_stride);
|
||||
|
||||
enqueue_kernel(ckPathTraceKernel, d_w, d_h);
|
||||
}
|
||||
|
||||
void thread_run(DeviceTask *task)
|
||||
{
|
||||
if(task->type == DeviceTask::FILM_CONVERT) {
|
||||
film_convert(*task, task->buffer, task->rgba_byte, task->rgba_half);
|
||||
}
|
||||
else if(task->type == DeviceTask::SHADER) {
|
||||
shader(*task);
|
||||
}
|
||||
else if(task->type == DeviceTask::RENDER) {
|
||||
RenderTile tile;
|
||||
DenoisingTask denoising(this, *task);
|
||||
|
||||
/* Keep rendering tiles until done. */
|
||||
while(task->acquire_tile(this, tile)) {
|
||||
if(tile.task == RenderTile::PATH_TRACE) {
|
||||
int start_sample = tile.start_sample;
|
||||
int end_sample = tile.start_sample + tile.num_samples;
|
||||
|
||||
for(int sample = start_sample; sample < end_sample; sample++) {
|
||||
if(task->get_cancel()) {
|
||||
if(task->need_finish_queue == false)
|
||||
break;
|
||||
}
|
||||
|
||||
path_trace(tile, sample);
|
||||
|
||||
tile.sample = sample + 1;
|
||||
|
||||
task->update_progress(&tile, tile.w*tile.h);
|
||||
}
|
||||
|
||||
/* Complete kernel execution before release tile */
|
||||
/* This helps in multi-device render;
|
||||
* The device that reaches the critical-section function
|
||||
* release_tile waits (stalling other devices from entering
|
||||
* release_tile) for all kernels to complete. If device1 (a
|
||||
* slow-render device) reaches release_tile first then it would
|
||||
* stall device2 (a fast-render device) from proceeding to render
|
||||
* next tile.
|
||||
*/
|
||||
clFinish(cqCommandQueue);
|
||||
}
|
||||
else if(tile.task == RenderTile::DENOISE) {
|
||||
tile.sample = tile.start_sample + tile.num_samples;
|
||||
denoise(tile, denoising);
|
||||
task->update_progress(&tile, tile.w*tile.h);
|
||||
}
|
||||
|
||||
task->release_tile(tile);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
bool is_split_kernel()
|
||||
{
|
||||
return false;
|
||||
}
|
||||
};
|
||||
|
||||
Device *opencl_create_mega_device(DeviceInfo& info, Stats& stats, Profiler &profiler, bool background)
|
||||
{
|
||||
return new OpenCLDeviceMegaKernel(info, stats, profiler, background);
|
||||
}
|
||||
|
||||
CCL_NAMESPACE_END
|
||||
|
||||
#endif
|
File diff suppressed because it is too large
Load Diff
|
@ -243,7 +243,7 @@ string OpenCLCache::get_kernel_md5()
|
|||
return self.kernel_md5;
|
||||
}
|
||||
|
||||
OpenCLDeviceBase::OpenCLProgram::OpenCLProgram(OpenCLDeviceBase *device,
|
||||
OpenCLDevice::OpenCLProgram::OpenCLProgram(OpenCLDevice *device,
|
||||
const string& program_name,
|
||||
const string& kernel_file,
|
||||
const string& kernel_build_options,
|
||||
|
@ -258,12 +258,12 @@ OpenCLDeviceBase::OpenCLProgram::OpenCLProgram(OpenCLDeviceBase *device,
|
|||
program = NULL;
|
||||
}
|
||||
|
||||
OpenCLDeviceBase::OpenCLProgram::~OpenCLProgram()
|
||||
OpenCLDevice::OpenCLProgram::~OpenCLProgram()
|
||||
{
|
||||
release();
|
||||
}
|
||||
|
||||
void OpenCLDeviceBase::OpenCLProgram::release()
|
||||
void OpenCLDevice::OpenCLProgram::release()
|
||||
{
|
||||
for(map<ustring, cl_kernel>::iterator kernel = kernels.begin(); kernel != kernels.end(); ++kernel) {
|
||||
if(kernel->second) {
|
||||
|
@ -277,7 +277,7 @@ void OpenCLDeviceBase::OpenCLProgram::release()
|
|||
}
|
||||
}
|
||||
|
||||
void OpenCLDeviceBase::OpenCLProgram::add_log(const string& msg, bool debug)
|
||||
void OpenCLDevice::OpenCLProgram::add_log(const string& msg, bool debug)
|
||||
{
|
||||
if(!use_stdout) {
|
||||
log += msg + "\n";
|
||||
|
@ -291,7 +291,7 @@ void OpenCLDeviceBase::OpenCLProgram::add_log(const string& msg, bool debug)
|
|||
}
|
||||
}
|
||||
|
||||
void OpenCLDeviceBase::OpenCLProgram::add_error(const string& msg)
|
||||
void OpenCLDevice::OpenCLProgram::add_error(const string& msg)
|
||||
{
|
||||
if(use_stdout) {
|
||||
fprintf(stderr, "%s\n", msg.c_str());
|
||||
|
@ -302,14 +302,14 @@ void OpenCLDeviceBase::OpenCLProgram::add_error(const string& msg)
|
|||
error_msg += msg;
|
||||
}
|
||||
|
||||
void OpenCLDeviceBase::OpenCLProgram::add_kernel(ustring name)
|
||||
void OpenCLDevice::OpenCLProgram::add_kernel(ustring name)
|
||||
{
|
||||
if(!kernels.count(name)) {
|
||||
kernels[name] = NULL;
|
||||
}
|
||||
}
|
||||
|
||||
bool OpenCLDeviceBase::OpenCLProgram::build_kernel(const string *debug_src)
|
||||
bool OpenCLDevice::OpenCLProgram::build_kernel(const string *debug_src)
|
||||
{
|
||||
string build_options;
|
||||
build_options = device->kernel_build_options(debug_src) + kernel_build_options;
|
||||
|
@ -341,7 +341,7 @@ bool OpenCLDeviceBase::OpenCLProgram::build_kernel(const string *debug_src)
|
|||
return (ciErr == CL_SUCCESS);
|
||||
}
|
||||
|
||||
bool OpenCLDeviceBase::OpenCLProgram::compile_kernel(const string *debug_src)
|
||||
bool OpenCLDevice::OpenCLProgram::compile_kernel(const string *debug_src)
|
||||
{
|
||||
string source = "#include \"kernel/kernels/opencl/" + kernel_file + "\"\n";
|
||||
/* We compile kernels consisting of many files. unfortunately OpenCL
|
||||
|
@ -389,14 +389,13 @@ static void escape_python_string(string& str)
|
|||
string_replace(str, "'", "\'");
|
||||
}
|
||||
|
||||
bool OpenCLDeviceBase::OpenCLProgram::compile_separate(const string& clbin)
|
||||
bool OpenCLDevice::OpenCLProgram::compile_separate(const string& clbin)
|
||||
{
|
||||
vector<string> args;
|
||||
args.push_back("--background");
|
||||
args.push_back("--factory-startup");
|
||||
args.push_back("--python-expr");
|
||||
|
||||
const char *force_all_platforms = (DebugFlags().opencl.kernel_type != DebugFlags::OpenCL::KERNEL_DEFAULT)? "true" : "false";
|
||||
int device_platform_id = device->device_num;
|
||||
string device_name = device->device_name;
|
||||
string platform_name = device->platform_name;
|
||||
|
@ -412,8 +411,7 @@ bool OpenCLDeviceBase::OpenCLProgram::compile_separate(const string& clbin)
|
|||
|
||||
args.push_back(
|
||||
string_printf(
|
||||
"import _cycles; _cycles.opencl_compile(r'%s', r'%d', r'%s', r'%s', r'%s', r'%s', r'%s')",
|
||||
force_all_platforms,
|
||||
"import _cycles; _cycles.opencl_compile(r'%d', r'%s', r'%s', r'%s', r'%s', r'%s')",
|
||||
device_platform_id,
|
||||
device_name.c_str(),
|
||||
platform_name.c_str(),
|
||||
|
@ -438,20 +436,19 @@ bool OpenCLDeviceBase::OpenCLProgram::compile_separate(const string& clbin)
|
|||
* module compile kernels. Parameters must match function above. */
|
||||
bool device_opencl_compile_kernel(const vector<string>& parameters)
|
||||
{
|
||||
bool force_all_platforms = parameters[0] == "true";
|
||||
int device_platform_id = std::stoi(parameters[1]);
|
||||
const string& device_name = parameters[2];
|
||||
const string& platform_name = parameters[3];
|
||||
const string& build_options = parameters[4];
|
||||
const string& kernel_file = parameters[5];
|
||||
const string& binary_path = parameters[6];
|
||||
int device_platform_id = std::stoi(parameters[0]);
|
||||
const string& device_name = parameters[1];
|
||||
const string& platform_name = parameters[2];
|
||||
const string& build_options = parameters[3];
|
||||
const string& kernel_file = parameters[4];
|
||||
const string& binary_path = parameters[5];
|
||||
|
||||
if(clewInit() != CLEW_SUCCESS) {
|
||||
return false;
|
||||
}
|
||||
|
||||
vector<OpenCLPlatformDevice> usable_devices;
|
||||
OpenCLInfo::get_usable_devices(&usable_devices, force_all_platforms);
|
||||
OpenCLInfo::get_usable_devices(&usable_devices);
|
||||
if(device_platform_id >= usable_devices.size()) {
|
||||
return false;
|
||||
}
|
||||
|
@ -504,7 +501,7 @@ bool device_opencl_compile_kernel(const vector<string>& parameters)
|
|||
return result;
|
||||
}
|
||||
|
||||
bool OpenCLDeviceBase::OpenCLProgram::load_binary(const string& clbin,
|
||||
bool OpenCLDevice::OpenCLProgram::load_binary(const string& clbin,
|
||||
const string *debug_src)
|
||||
{
|
||||
/* read binary into memory */
|
||||
|
@ -535,7 +532,7 @@ bool OpenCLDeviceBase::OpenCLProgram::load_binary(const string& clbin,
|
|||
return true;
|
||||
}
|
||||
|
||||
bool OpenCLDeviceBase::OpenCLProgram::save_binary(const string& clbin)
|
||||
bool OpenCLDevice::OpenCLProgram::save_binary(const string& clbin)
|
||||
{
|
||||
size_t size = 0;
|
||||
clGetProgramInfo(program, CL_PROGRAM_BINARY_SIZES, sizeof(size_t), &size, NULL);
|
||||
|
@ -551,7 +548,7 @@ bool OpenCLDeviceBase::OpenCLProgram::save_binary(const string& clbin)
|
|||
return path_write_binary(clbin, binary);
|
||||
}
|
||||
|
||||
void OpenCLDeviceBase::OpenCLProgram::load()
|
||||
void OpenCLDevice::OpenCLProgram::load()
|
||||
{
|
||||
assert(device);
|
||||
|
||||
|
@ -642,7 +639,7 @@ void OpenCLDeviceBase::OpenCLProgram::load()
|
|||
loaded = true;
|
||||
}
|
||||
|
||||
void OpenCLDeviceBase::OpenCLProgram::report_error()
|
||||
void OpenCLDevice::OpenCLProgram::report_error()
|
||||
{
|
||||
/* If loaded is true, there was no error. */
|
||||
if(loaded) return;
|
||||
|
@ -656,13 +653,13 @@ void OpenCLDeviceBase::OpenCLProgram::report_error()
|
|||
}
|
||||
}
|
||||
|
||||
cl_kernel OpenCLDeviceBase::OpenCLProgram::operator()()
|
||||
cl_kernel OpenCLDevice::OpenCLProgram::operator()()
|
||||
{
|
||||
assert(kernels.size() == 1);
|
||||
return kernels.begin()->second;
|
||||
}
|
||||
|
||||
cl_kernel OpenCLDeviceBase::OpenCLProgram::operator()(ustring name)
|
||||
cl_kernel OpenCLDevice::OpenCLProgram::operator()(ustring name)
|
||||
{
|
||||
assert(kernels.count(name));
|
||||
return kernels[name];
|
||||
|
@ -716,28 +713,6 @@ bool OpenCLInfo::kernel_use_advanced_shading(const string& platform)
|
|||
return false;
|
||||
}
|
||||
|
||||
bool OpenCLInfo::kernel_use_split(const string& platform_name,
|
||||
const cl_device_type device_type)
|
||||
{
|
||||
if(DebugFlags().opencl.kernel_type == DebugFlags::OpenCL::KERNEL_SPLIT) {
|
||||
VLOG(1) << "Forcing split kernel to use.";
|
||||
return true;
|
||||
}
|
||||
if(DebugFlags().opencl.kernel_type == DebugFlags::OpenCL::KERNEL_MEGA) {
|
||||
VLOG(1) << "Forcing mega kernel to use.";
|
||||
return false;
|
||||
}
|
||||
/* TODO(sergey): Replace string lookups with more enum-like API,
|
||||
* similar to device/vendor checks blender's gpu.
|
||||
*/
|
||||
if(platform_name == "AMD Accelerated Parallel Processing" &&
|
||||
device_type == CL_DEVICE_TYPE_GPU)
|
||||
{
|
||||
return true;
|
||||
}
|
||||
return false;
|
||||
}
|
||||
|
||||
bool OpenCLInfo::device_supported(const string& platform_name,
|
||||
const cl_device_id device_id)
|
||||
{
|
||||
|
@ -878,8 +853,6 @@ string OpenCLInfo::get_hardware_id(const string& platform_name, cl_device_id dev
|
|||
void OpenCLInfo::get_usable_devices(vector<OpenCLPlatformDevice> *usable_devices,
|
||||
bool force_all)
|
||||
{
|
||||
const bool force_all_platforms = force_all ||
|
||||
(DebugFlags().opencl.kernel_type != DebugFlags::OpenCL::KERNEL_DEFAULT);
|
||||
const cl_device_type device_type = OpenCLInfo::device_type();
|
||||
static bool first_time = true;
|
||||
#define FIRST_VLOG(severity) if(first_time) VLOG(severity)
|
||||
|
@ -952,7 +925,7 @@ void OpenCLInfo::get_usable_devices(vector<OpenCLPlatformDevice> *usable_devices
|
|||
<< " due to old compiler version.";
|
||||
continue;
|
||||
}
|
||||
if(force_all_platforms ||
|
||||
if(force_all ||
|
||||
device_supported(platform_name, device_id))
|
||||
{
|
||||
cl_device_type device_type;
|
||||
|
|
|
@ -36,8 +36,8 @@ set(SRC_CUDA_KERNELS
|
|||
)
|
||||
|
||||
set(SRC_OPENCL_KERNELS
|
||||
kernels/opencl/kernel.cl
|
||||
kernels/opencl/kernel_bake.cl
|
||||
kernels/opencl/kernel_base.cl
|
||||
kernels/opencl/kernel_displace.cl
|
||||
kernels/opencl/kernel_background.cl
|
||||
kernels/opencl/kernel_state_buffer_size.cl
|
||||
|
|
|
@ -14,72 +14,14 @@
|
|||
* limitations under the License.
|
||||
*/
|
||||
|
||||
/* OpenCL kernel entry points - unfinished */
|
||||
/* OpenCL base kernels entry points */
|
||||
|
||||
#include "kernel/kernel_compat_opencl.h"
|
||||
#include "kernel/kernel_math.h"
|
||||
#include "kernel/kernel_types.h"
|
||||
#include "kernel/kernel_globals.h"
|
||||
#include "kernel/kernel_color.h"
|
||||
#include "kernel/kernels/opencl/kernel_opencl_image.h"
|
||||
|
||||
#include "kernel/kernel_film.h"
|
||||
|
||||
#if defined(__COMPILE_ONLY_MEGAKERNEL__) || !defined(__NO_BAKING__)
|
||||
# include "kernel/kernel_path.h"
|
||||
# include "kernel/kernel_path_branched.h"
|
||||
#else /* __COMPILE_ONLY_MEGAKERNEL__ */
|
||||
/* Include only actually used headers for the case
|
||||
* when path tracing kernels are not needed.
|
||||
*/
|
||||
# include "kernel/kernel_random.h"
|
||||
# include "kernel/kernel_differential.h"
|
||||
# include "kernel/kernel_montecarlo.h"
|
||||
# include "kernel/kernel_projection.h"
|
||||
# include "kernel/geom/geom.h"
|
||||
# include "kernel/bvh/bvh.h"
|
||||
|
||||
# include "kernel/kernel_accumulate.h"
|
||||
# include "kernel/kernel_camera.h"
|
||||
# include "kernel/kernel_shader.h"
|
||||
#endif /* defined(__COMPILE_ONLY_MEGAKERNEL__) || !defined(__NO_BAKING__) */
|
||||
|
||||
#include "kernel/kernel_bake.h"
|
||||
|
||||
#ifdef __COMPILE_ONLY_MEGAKERNEL__
|
||||
|
||||
__kernel void kernel_ocl_path_trace(
|
||||
ccl_constant KernelData *data,
|
||||
ccl_global float *buffer,
|
||||
|
||||
KERNEL_BUFFER_PARAMS,
|
||||
|
||||
int sample,
|
||||
int sx, int sy, int sw, int sh, int offset, int stride)
|
||||
{
|
||||
KernelGlobals kglobals, *kg = &kglobals;
|
||||
|
||||
kg->data = data;
|
||||
|
||||
kernel_set_buffer_pointers(kg, KERNEL_BUFFER_ARGS);
|
||||
kernel_set_buffer_info(kg);
|
||||
|
||||
int x = sx + ccl_global_id(0);
|
||||
int y = sy + ccl_global_id(1);
|
||||
bool thread_is_active = x < sx + sw && y < sy + sh;
|
||||
if(thread_is_active) {
|
||||
kernel_path_trace(kg, buffer, sample, x, y, offset, stride);
|
||||
}
|
||||
if(kernel_data.film.cryptomatte_passes) {
|
||||
/* Make sure no thread is writing to the buffers. */
|
||||
ccl_barrier(CCL_LOCAL_MEM_FENCE);
|
||||
if(thread_is_active) {
|
||||
kernel_cryptomatte_post(kg, buffer, sample, x, y, offset, stride);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
#else /* __COMPILE_ONLY_MEGAKERNEL__ */
|
||||
|
||||
__kernel void kernel_ocl_convert_to_byte(
|
||||
ccl_constant KernelData *data,
|
||||
|
@ -144,5 +86,3 @@ __kernel void kernel_ocl_zero_buffer(ccl_global float4 *buffer, uint64_t size, u
|
|||
}
|
||||
}
|
||||
}
|
||||
|
||||
#endif /* __COMPILE_ONLY_MEGAKERNEL__ */
|
|
@ -90,7 +90,6 @@ void DebugFlags::CUDA::reset()
|
|||
|
||||
DebugFlags::OpenCL::OpenCL()
|
||||
: device_type(DebugFlags::OpenCL::DEVICE_ALL),
|
||||
kernel_type(DebugFlags::OpenCL::KERNEL_DEFAULT),
|
||||
debug(false),
|
||||
single_program(false)
|
||||
{
|
||||
|
@ -122,14 +121,6 @@ void DebugFlags::OpenCL::reset()
|
|||
device_type = DebugFlags::OpenCL::DEVICE_ACCELERATOR;
|
||||
}
|
||||
}
|
||||
/* Initialize kernel type from environment variables. */
|
||||
kernel_type = DebugFlags::OpenCL::KERNEL_DEFAULT;
|
||||
if(getenv("CYCLES_OPENCL_MEGA_KERNEL_TEST") != NULL) {
|
||||
kernel_type = DebugFlags::OpenCL::KERNEL_MEGA;
|
||||
}
|
||||
else if(getenv("CYCLES_OPENCL_SPLIT_KERNEL_TEST") != NULL) {
|
||||
kernel_type = DebugFlags::OpenCL::KERNEL_SPLIT;
|
||||
}
|
||||
/* Initialize other flags from environment variables. */
|
||||
debug = (getenv("CYCLES_OPENCL_DEBUG") != NULL);
|
||||
single_program = (getenv("CYCLES_OPENCL_SINGLE_PROGRAM") != NULL);
|
||||
|
@ -164,8 +155,7 @@ std::ostream& operator <<(std::ostream &os,
|
|||
os << "CUDA flags:\n"
|
||||
<< " Adaptive Compile: " << string_from_bool(debug_flags.cuda.adaptive_compile) << "\n";
|
||||
|
||||
const char *opencl_device_type,
|
||||
*opencl_kernel_type;
|
||||
const char *opencl_device_type;
|
||||
switch(debug_flags.opencl.device_type) {
|
||||
case DebugFlags::OpenCL::DEVICE_NONE:
|
||||
opencl_device_type = "NONE";
|
||||
|
@ -186,20 +176,8 @@ std::ostream& operator <<(std::ostream &os,
|
|||
opencl_device_type = "ACCELERATOR";
|
||||
break;
|
||||
}
|
||||
switch(debug_flags.opencl.kernel_type) {
|
||||
case DebugFlags::OpenCL::KERNEL_DEFAULT:
|
||||
opencl_kernel_type = "DEFAULT";
|
||||
break;
|
||||
case DebugFlags::OpenCL::KERNEL_MEGA:
|
||||
opencl_kernel_type = "MEGA";
|
||||
break;
|
||||
case DebugFlags::OpenCL::KERNEL_SPLIT:
|
||||
opencl_kernel_type = "SPLIT";
|
||||
break;
|
||||
}
|
||||
os << "OpenCL flags:\n"
|
||||
<< " Device type : " << opencl_device_type << "\n"
|
||||
<< " Kernel type : " << opencl_kernel_type << "\n"
|
||||
<< " Debug : " << string_from_bool(debug_flags.opencl.debug) << "\n"
|
||||
<< " Single program : " << string_from_bool(debug_flags.opencl.single_program) << "\n"
|
||||
<< " Memory limit : " << string_human_readable_size(debug_flags.opencl.mem_limit) << "\n";
|
||||
|
|
|
@ -123,9 +123,6 @@ public:
|
|||
/* Requested device type. */
|
||||
DeviceType device_type;
|
||||
|
||||
/* Requested kernel type. */
|
||||
KernelType kernel_type;
|
||||
|
||||
/* Use debug version of the kernel. */
|
||||
bool debug;
|
||||
|
||||
|
|
Loading…
Reference in New Issue