Cr8-xyz/blender
1
0
Fork 0
Code Issues Pull Requests Actions 1 Packages Projects Releases Wiki Activity
Files
a0c02e4d1b87f4e83c0dfe794f28482030be9896
blender/intern/cycles/device/opencl/opencl_split.cpp

460 lines
15 KiB
C++
Raw Normal View History

Cycles: Split device_opencl.cpp into multiple files for easier maintenance There are no user-visible changes, just some internal restructuring. Differential Revision: https://developer.blender.org/D2231
2016-09-14 23:47:54 +02:00
/*
* 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
Cycles: Make all #include statements relative to cycles source directory The idea is to make include statements more explicit and obvious where the file is coming from, additionally reducing chance of wrong header being picked up. For example, it was not obvious whether bvh.h was refferring to builder or traversal, whenter node.h is a generic graph node or a shader node and cases like that. Surely this might look obvious for the active developers, but after some time of not touching the code it becomes less obvious where file is coming from. This was briefly mentioned in T50824 and seems @brecht is fine with such explicitness, but need to agree with all active developers before committing this. Please note that this patch is lacking changes related on GPU/OpenCL support. This will be solved if/when we all agree this is a good idea to move forward. Reviewers: brecht, lukasstockner97, maiself, nirved, dingto, juicyfruit, swerner Reviewed By: lukasstockner97, maiself, nirved, dingto Subscribers: brecht Differential Revision: https://developer.blender.org/D2586
2017-03-28 20:39:14 +02:00
#include "device/opencl/opencl.h"
Cycles: Split device_opencl.cpp into multiple files for easier maintenance There are no user-visible changes, just some internal restructuring. Differential Revision: https://developer.blender.org/D2231
2016-09-14 23:47:54 +02:00
Cycles: Make all #include statements relative to cycles source directory The idea is to make include statements more explicit and obvious where the file is coming from, additionally reducing chance of wrong header being picked up. For example, it was not obvious whether bvh.h was refferring to builder or traversal, whenter node.h is a generic graph node or a shader node and cases like that. Surely this might look obvious for the active developers, but after some time of not touching the code it becomes less obvious where file is coming from. This was briefly mentioned in T50824 and seems @brecht is fine with such explicitness, but need to agree with all active developers before committing this. Please note that this patch is lacking changes related on GPU/OpenCL support. This will be solved if/when we all agree this is a good idea to move forward. Reviewers: brecht, lukasstockner97, maiself, nirved, dingto, juicyfruit, swerner Reviewed By: lukasstockner97, maiself, nirved, dingto Subscribers: brecht Differential Revision: https://developer.blender.org/D2586
2017-03-28 20:39:14 +02:00
#include "render/buffers.h"
Cycles: Split device_opencl.cpp into multiple files for easier maintenance There are no user-visible changes, just some internal restructuring. Differential Revision: https://developer.blender.org/D2231
2016-09-14 23:47:54 +02:00
Cycles: Make all #include statements relative to cycles source directory The idea is to make include statements more explicit and obvious where the file is coming from, additionally reducing chance of wrong header being picked up. For example, it was not obvious whether bvh.h was refferring to builder or traversal, whenter node.h is a generic graph node or a shader node and cases like that. Surely this might look obvious for the active developers, but after some time of not touching the code it becomes less obvious where file is coming from. This was briefly mentioned in T50824 and seems @brecht is fine with such explicitness, but need to agree with all active developers before committing this. Please note that this patch is lacking changes related on GPU/OpenCL support. This will be solved if/when we all agree this is a good idea to move forward. Reviewers: brecht, lukasstockner97, maiself, nirved, dingto, juicyfruit, swerner Reviewed By: lukasstockner97, maiself, nirved, dingto Subscribers: brecht Differential Revision: https://developer.blender.org/D2586
2017-03-28 20:39:14 +02:00
#include "kernel/kernel_types.h"
#include "kernel/split/kernel_split_data_types.h"
Cycles: Split device_opencl.cpp into multiple files for easier maintenance There are no user-visible changes, just some internal restructuring. Differential Revision: https://developer.blender.org/D2231
2016-09-14 23:47:54 +02:00
Cycles: Make all #include statements relative to cycles source directory The idea is to make include statements more explicit and obvious where the file is coming from, additionally reducing chance of wrong header being picked up. For example, it was not obvious whether bvh.h was refferring to builder or traversal, whenter node.h is a generic graph node or a shader node and cases like that. Surely this might look obvious for the active developers, but after some time of not touching the code it becomes less obvious where file is coming from. This was briefly mentioned in T50824 and seems @brecht is fine with such explicitness, but need to agree with all active developers before committing this. Please note that this patch is lacking changes related on GPU/OpenCL support. This will be solved if/when we all agree this is a good idea to move forward. Reviewers: brecht, lukasstockner97, maiself, nirved, dingto, juicyfruit, swerner Reviewed By: lukasstockner97, maiself, nirved, dingto Subscribers: brecht Differential Revision: https://developer.blender.org/D2586
2017-03-28 20:39:14 +02:00
#include "device/device_split_kernel.h"
Cycles: OpenCL split kernel refactor This does a few things at once: - Refactors host side split kernel logic into a new device agnostic class `DeviceSplitKernel`. - Removes tile splitting, a new work pool implementation takes its place and allows as many threads as will fit in memory regardless of tile size, which can give performance gains. - Refactors split state buffers into one buffer, as well as reduces the number of arguments passed to kernels. Means there's less code to deal with overall. - Moves kernel logic out of OpenCL kernel files so they can later be used by other device types. - Replaced OpenCL specific APIs with new generic versions - Tiles can now be seen updating during rendering
2017-02-22 08:10:02 -05:00
Cycles: Add artificial memory limit debug option for OpenCL
2017-07-05 20:16:41 -04:00
#include "util/util_algorithm.h"
Cycles: Make all #include statements relative to cycles source directory The idea is to make include statements more explicit and obvious where the file is coming from, additionally reducing chance of wrong header being picked up. For example, it was not obvious whether bvh.h was refferring to builder or traversal, whenter node.h is a generic graph node or a shader node and cases like that. Surely this might look obvious for the active developers, but after some time of not touching the code it becomes less obvious where file is coming from. This was briefly mentioned in T50824 and seems @brecht is fine with such explicitness, but need to agree with all active developers before committing this. Please note that this patch is lacking changes related on GPU/OpenCL support. This will be solved if/when we all agree this is a good idea to move forward. Reviewers: brecht, lukasstockner97, maiself, nirved, dingto, juicyfruit, swerner Reviewed By: lukasstockner97, maiself, nirved, dingto Subscribers: brecht Differential Revision: https://developer.blender.org/D2586
2017-03-28 20:39:14 +02:00
#include "util/util_logging.h"
#include "util/util_md5.h"
#include "util/util_path.h"
#include "util/util_time.h"
Cycles: Split device_opencl.cpp into multiple files for easier maintenance There are no user-visible changes, just some internal restructuring. Differential Revision: https://developer.blender.org/D2231
2016-09-14 23:47:54 +02:00
CCL_NAMESPACE_BEGIN
Cycles: OpenCL split kernel refactor This does a few things at once: - Refactors host side split kernel logic into a new device agnostic class `DeviceSplitKernel`. - Removes tile splitting, a new work pool implementation takes its place and allows as many threads as will fit in memory regardless of tile size, which can give performance gains. - Refactors split state buffers into one buffer, as well as reduces the number of arguments passed to kernels. Means there's less code to deal with overall. - Moves kernel logic out of OpenCL kernel files so they can later be used by other device types. - Replaced OpenCL specific APIs with new generic versions - Tiles can now be seen updating during rendering
2017-02-22 08:10:02 -05:00
class OpenCLSplitKernel;
Cycles: Split device_opencl.cpp into multiple files for easier maintenance There are no user-visible changes, just some internal restructuring. Differential Revision: https://developer.blender.org/D2231
2016-09-14 23:47:54 +02:00
Cycles: Fix compilation error without C++11
2017-10-24 11:14:01 +02:00
namespace {
/* Copy dummy KernelGlobals related to OpenCL from kernel_globals.h to
* fetch its size.
*/
typedef struct KernelGlobalsDummy {
ccl_constant KernelData *data;
ccl_global char *buffers[8];
#define KERNEL_TEX(type, name) \
TextureInfo name;
# include "kernel/kernel_textures.h"
#undef KERNEL_TEX
SplitData split_data;
SplitParams split_param_data;
} KernelGlobalsDummy;
} // namespace
Cycles: OpenCL split kernel refactor This does a few things at once: - Refactors host side split kernel logic into a new device agnostic class `DeviceSplitKernel`. - Removes tile splitting, a new work pool implementation takes its place and allows as many threads as will fit in memory regardless of tile size, which can give performance gains. - Refactors split state buffers into one buffer, as well as reduces the number of arguments passed to kernels. Means there's less code to deal with overall. - Moves kernel logic out of OpenCL kernel files so they can later be used by other device types. - Replaced OpenCL specific APIs with new generic versions - Tiles can now be seen updating during rendering
2017-02-22 08:10:02 -05:00
static string get_build_options(OpenCLDeviceBase *device, const DeviceRequestedFeatures& requested_features)
{
string build_options = "-D__SPLIT_KERNEL__ ";
build_options += requested_features.get_build_options();
/* Set compute device build option. */
cl_device_type device_type;
Cycles: Simplify some extra OpenCL query code
2017-03-20 18:59:46 +01:00
OpenCLInfo::get_device_type(device->cdDevice, &device_type, &device->ciErr);
Cycles: OpenCL split kernel refactor This does a few things at once: - Refactors host side split kernel logic into a new device agnostic class `DeviceSplitKernel`. - Removes tile splitting, a new work pool implementation takes its place and allows as many threads as will fit in memory regardless of tile size, which can give performance gains. - Refactors split state buffers into one buffer, as well as reduces the number of arguments passed to kernels. Means there's less code to deal with overall. - Moves kernel logic out of OpenCL kernel files so they can later be used by other device types. - Replaced OpenCL specific APIs with new generic versions - Tiles can now be seen updating during rendering
2017-02-22 08:10:02 -05:00
assert(device->ciErr == CL_SUCCESS);
if(device_type == CL_DEVICE_TYPE_GPU) {
build_options += " -D__COMPUTE_DEVICE_GPU__";
Cycles: Split device_opencl.cpp into multiple files for easier maintenance There are no user-visible changes, just some internal restructuring. Differential Revision: https://developer.blender.org/D2231
2016-09-14 23:47:54 +02:00
}
Cycles: OpenCL split kernel refactor This does a few things at once: - Refactors host side split kernel logic into a new device agnostic class `DeviceSplitKernel`. - Removes tile splitting, a new work pool implementation takes its place and allows as many threads as will fit in memory regardless of tile size, which can give performance gains. - Refactors split state buffers into one buffer, as well as reduces the number of arguments passed to kernels. Means there's less code to deal with overall. - Moves kernel logic out of OpenCL kernel files so they can later be used by other device types. - Replaced OpenCL specific APIs with new generic versions - Tiles can now be seen updating during rendering
2017-02-22 08:10:02 -05:00
return build_options;
}
Cycles: Split device_opencl.cpp into multiple files for easier maintenance There are no user-visible changes, just some internal restructuring. Differential Revision: https://developer.blender.org/D2231
2016-09-14 23:47:54 +02:00
/* OpenCLDeviceSplitKernel's declaration/definition. */
class OpenCLDeviceSplitKernel : public OpenCLDeviceBase
{
public:
Cycles: OpenCL split kernel refactor This does a few things at once: - Refactors host side split kernel logic into a new device agnostic class `DeviceSplitKernel`. - Removes tile splitting, a new work pool implementation takes its place and allows as many threads as will fit in memory regardless of tile size, which can give performance gains. - Refactors split state buffers into one buffer, as well as reduces the number of arguments passed to kernels. Means there's less code to deal with overall. - Moves kernel logic out of OpenCL kernel files so they can later be used by other device types. - Replaced OpenCL specific APIs with new generic versions - Tiles can now be seen updating during rendering
2017-02-22 08:10:02 -05:00
DeviceSplitKernel *split_kernel;
Cycles: Split device_opencl.cpp into multiple files for easier maintenance There are no user-visible changes, just some internal restructuring. Differential Revision: https://developer.blender.org/D2231
2016-09-14 23:47:54 +02:00
OpenCLProgram program_data_init;
Cycles: Calculate size of split state buffer kernel side By calculating the size of the state buffer in the kernel rather than the host less code is needed and the size actually reflects the requested features. Will also be a little faster in some cases because of larger global work size.
2017-03-04 06:29:01 -05:00
OpenCLProgram program_state_buffer_size;
Cycles: Split device_opencl.cpp into multiple files for easier maintenance There are no user-visible changes, just some internal restructuring. Differential Revision: https://developer.blender.org/D2231
2016-09-14 23:47:54 +02:00
Cycles: OpenCL split kernel refactor This does a few things at once: - Refactors host side split kernel logic into a new device agnostic class `DeviceSplitKernel`. - Removes tile splitting, a new work pool implementation takes its place and allows as many threads as will fit in memory regardless of tile size, which can give performance gains. - Refactors split state buffers into one buffer, as well as reduces the number of arguments passed to kernels. Means there's less code to deal with overall. - Moves kernel logic out of OpenCL kernel files so they can later be used by other device types. - Replaced OpenCL specific APIs with new generic versions - Tiles can now be seen updating during rendering
2017-02-22 08:10:02 -05:00
OpenCLDeviceSplitKernel(DeviceInfo& info, Stats &stats, bool background_);
Cycles: Split device_opencl.cpp into multiple files for easier maintenance There are no user-visible changes, just some internal restructuring. Differential Revision: https://developer.blender.org/D2231
2016-09-14 23:47:54 +02:00
Cycles: OpenCL split kernel refactor This does a few things at once: - Refactors host side split kernel logic into a new device agnostic class `DeviceSplitKernel`. - Removes tile splitting, a new work pool implementation takes its place and allows as many threads as will fit in memory regardless of tile size, which can give performance gains. - Refactors split state buffers into one buffer, as well as reduces the number of arguments passed to kernels. Means there's less code to deal with overall. - Moves kernel logic out of OpenCL kernel files so they can later be used by other device types. - Replaced OpenCL specific APIs with new generic versions - Tiles can now be seen updating during rendering
2017-02-22 08:10:02 -05:00
~OpenCLDeviceSplitKernel()
{
task_pool.stop();
Cycles: Split device_opencl.cpp into multiple files for easier maintenance There are no user-visible changes, just some internal restructuring. Differential Revision: https://developer.blender.org/D2231
2016-09-14 23:47:54 +02:00
Cycles: OpenCL split kernel refactor This does a few things at once: - Refactors host side split kernel logic into a new device agnostic class `DeviceSplitKernel`. - Removes tile splitting, a new work pool implementation takes its place and allows as many threads as will fit in memory regardless of tile size, which can give performance gains. - Refactors split state buffers into one buffer, as well as reduces the number of arguments passed to kernels. Means there's less code to deal with overall. - Moves kernel logic out of OpenCL kernel files so they can later be used by other device types. - Replaced OpenCL specific APIs with new generic versions - Tiles can now be seen updating during rendering
2017-02-22 08:10:02 -05:00
/* Release kernels */
program_data_init.release();
Cycles: Split device_opencl.cpp into multiple files for easier maintenance There are no user-visible changes, just some internal restructuring. Differential Revision: https://developer.blender.org/D2231
2016-09-14 23:47:54 +02:00
Cycles: OpenCL split kernel refactor This does a few things at once: - Refactors host side split kernel logic into a new device agnostic class `DeviceSplitKernel`. - Removes tile splitting, a new work pool implementation takes its place and allows as many threads as will fit in memory regardless of tile size, which can give performance gains. - Refactors split state buffers into one buffer, as well as reduces the number of arguments passed to kernels. Means there's less code to deal with overall. - Moves kernel logic out of OpenCL kernel files so they can later be used by other device types. - Replaced OpenCL specific APIs with new generic versions - Tiles can now be seen updating during rendering
2017-02-22 08:10:02 -05:00
delete split_kernel;
}
Cycles: Split device_opencl.cpp into multiple files for easier maintenance There are no user-visible changes, just some internal restructuring. Differential Revision: https://developer.blender.org/D2231
2016-09-14 23:47:54 +02:00
Cycles: Show samples progress for OpenCL split kernel
2017-05-05 13:31:21 +02:00
virtual bool show_samples() const {
return true;
}
Cycles: OpenCL split kernel refactor This does a few things at once: - Refactors host side split kernel logic into a new device agnostic class `DeviceSplitKernel`. - Removes tile splitting, a new work pool implementation takes its place and allows as many threads as will fit in memory regardless of tile size, which can give performance gains. - Refactors split state buffers into one buffer, as well as reduces the number of arguments passed to kernels. Means there's less code to deal with overall. - Moves kernel logic out of OpenCL kernel files so they can later be used by other device types. - Replaced OpenCL specific APIs with new generic versions - Tiles can now be seen updating during rendering
2017-02-22 08:10:02 -05:00
virtual bool load_kernels(const DeviceRequestedFeatures& requested_features,
vector<OpenCLDeviceBase::OpenCLProgram*> &programs)
Cycles: Split device_opencl.cpp into multiple files for easier maintenance There are no user-visible changes, just some internal restructuring. Differential Revision: https://developer.blender.org/D2231
2016-09-14 23:47:54 +02:00
{
Cycles: add single program debug option for split kernel Single program generally compiles kernels faster (2-3 times), loads faster, takes less drive space (2-3 times), and reduces the number of cached kernels.
2017-03-08 17:56:06 +01:00
bool single_program = OpenCLInfo::use_single_program();
Cycles: OpenCL split kernel refactor This does a few things at once: - Refactors host side split kernel logic into a new device agnostic class `DeviceSplitKernel`. - Removes tile splitting, a new work pool implementation takes its place and allows as many threads as will fit in memory regardless of tile size, which can give performance gains. - Refactors split state buffers into one buffer, as well as reduces the number of arguments passed to kernels. Means there's less code to deal with overall. - Moves kernel logic out of OpenCL kernel files so they can later be used by other device types. - Replaced OpenCL specific APIs with new generic versions - Tiles can now be seen updating during rendering
2017-02-22 08:10:02 -05:00
program_data_init = OpenCLDeviceBase::OpenCLProgram(this,
Cycles: add single program debug option for split kernel Single program generally compiles kernels faster (2-3 times), loads faster, takes less drive space (2-3 times), and reduces the number of cached kernels.
2017-03-08 17:56:06 +01:00
single_program ? "split" : "split_data_init",
single_program ? "kernel_split.cl" : "kernel_data_init.cl",
Cycles: OpenCL split kernel refactor This does a few things at once: - Refactors host side split kernel logic into a new device agnostic class `DeviceSplitKernel`. - Removes tile splitting, a new work pool implementation takes its place and allows as many threads as will fit in memory regardless of tile size, which can give performance gains. - Refactors split state buffers into one buffer, as well as reduces the number of arguments passed to kernels. Means there's less code to deal with overall. - Moves kernel logic out of OpenCL kernel files so they can later be used by other device types. - Replaced OpenCL specific APIs with new generic versions - Tiles can now be seen updating during rendering
2017-02-22 08:10:02 -05:00
get_build_options(this, requested_features));
Cycles: add single program debug option for split kernel Single program generally compiles kernels faster (2-3 times), loads faster, takes less drive space (2-3 times), and reduces the number of cached kernels.
2017-03-08 17:56:06 +01:00
Cycles: OpenCL split kernel refactor This does a few things at once: - Refactors host side split kernel logic into a new device agnostic class `DeviceSplitKernel`. - Removes tile splitting, a new work pool implementation takes its place and allows as many threads as will fit in memory regardless of tile size, which can give performance gains. - Refactors split state buffers into one buffer, as well as reduces the number of arguments passed to kernels. Means there's less code to deal with overall. - Moves kernel logic out of OpenCL kernel files so they can later be used by other device types. - Replaced OpenCL specific APIs with new generic versions - Tiles can now be seen updating during rendering
2017-02-22 08:10:02 -05:00
program_data_init.add_kernel(ustring("path_trace_data_init"));
programs.push_back(&program_data_init);
Cycles: Calculate size of split state buffer kernel side By calculating the size of the state buffer in the kernel rather than the host less code is needed and the size actually reflects the requested features. Will also be a little faster in some cases because of larger global work size.
2017-03-04 06:29:01 -05:00
program_state_buffer_size = OpenCLDeviceBase::OpenCLProgram(this,
Cycles: add single program debug option for split kernel Single program generally compiles kernels faster (2-3 times), loads faster, takes less drive space (2-3 times), and reduces the number of cached kernels.
2017-03-08 17:56:06 +01:00
single_program ? "split" : "split_state_buffer_size",
single_program ? "kernel_split.cl" : "kernel_state_buffer_size.cl",
Cycles: Calculate size of split state buffer kernel side By calculating the size of the state buffer in the kernel rather than the host less code is needed and the size actually reflects the requested features. Will also be a little faster in some cases because of larger global work size.
2017-03-04 06:29:01 -05:00
get_build_options(this, requested_features));
program_state_buffer_size.add_kernel(ustring("path_trace_state_buffer_size"));
programs.push_back(&program_state_buffer_size);
Cycles: OpenCL split kernel refactor This does a few things at once: - Refactors host side split kernel logic into a new device agnostic class `DeviceSplitKernel`. - Removes tile splitting, a new work pool implementation takes its place and allows as many threads as will fit in memory regardless of tile size, which can give performance gains. - Refactors split state buffers into one buffer, as well as reduces the number of arguments passed to kernels. Means there's less code to deal with overall. - Moves kernel logic out of OpenCL kernel files so they can later be used by other device types. - Replaced OpenCL specific APIs with new generic versions - Tiles can now be seen updating during rendering
2017-02-22 08:10:02 -05:00
return split_kernel->load_kernels(requested_features);
}
Cycles: Split device_opencl.cpp into multiple files for easier maintenance There are no user-visible changes, just some internal restructuring. Differential Revision: https://developer.blender.org/D2231
2016-09-14 23:47:54 +02:00
Cycles: OpenCL split kernel refactor This does a few things at once: - Refactors host side split kernel logic into a new device agnostic class `DeviceSplitKernel`. - Removes tile splitting, a new work pool implementation takes its place and allows as many threads as will fit in memory regardless of tile size, which can give performance gains. - Refactors split state buffers into one buffer, as well as reduces the number of arguments passed to kernels. Means there's less code to deal with overall. - Moves kernel logic out of OpenCL kernel files so they can later be used by other device types. - Replaced OpenCL specific APIs with new generic versions - Tiles can now be seen updating during rendering
2017-02-22 08:10:02 -05:00
void thread_run(DeviceTask *task)
{
Cycles: Pack kernel textures into buffers for OpenCL Image textures were being packed into a single buffer for OpenCL, which limited the amount of memory available for images to the size of one buffer (usually 4gb on AMD hardware). By packing textures into multiple buffers that limit is removed, while simultaneously reducing the number of buffers that need to be passed to each kernel. Benchmarks were within 2%. Fixes T51554. Differential Revision: https://developer.blender.org/D2745
2017-08-08 07:12:04 -04:00
flush_texture_buffers();
Cycles: OpenCL split kernel refactor This does a few things at once: - Refactors host side split kernel logic into a new device agnostic class `DeviceSplitKernel`. - Removes tile splitting, a new work pool implementation takes its place and allows as many threads as will fit in memory regardless of tile size, which can give performance gains. - Refactors split state buffers into one buffer, as well as reduces the number of arguments passed to kernels. Means there's less code to deal with overall. - Moves kernel logic out of OpenCL kernel files so they can later be used by other device types. - Replaced OpenCL specific APIs with new generic versions - Tiles can now be seen updating during rendering
2017-02-22 08:10:02 -05:00
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);
}
Cycles: Implement denoising option for reducing noise in the rendered image This commit contains the first part of the new Cycles denoising option, which filters the resulting image using information gathered during rendering to get rid of noise while preserving visual features as well as possible. To use the option, enable it in the render layer options. The default settings fit a wide range of scenes, but the user can tweak individual settings to control the tradeoff between a noise-free image, image details, and calculation time. Note that the denoiser may still change in the future and that some features are not implemented yet. The most important missing feature is animation denoising, which uses information from multiple frames at once to produce a flicker-free and smoother result. These features will be added in the future. Finally, thanks to all the people who supported this project: - Google (through the GSoC) and Theory Studios for sponsoring the development - The authors of the papers I used for implementing the denoiser (more details on them will be included in the technical docs) - The other Cycles devs for feedback on the code, especially Sergey for mentoring the GSoC project and Brecht for the code review! - And of course the users who helped with testing, reported bugs and things that could and/or should work better!
2017-05-07 14:40:58 +02:00
else if(task->type == DeviceTask::RENDER) {
Cycles: OpenCL split kernel refactor This does a few things at once: - Refactors host side split kernel logic into a new device agnostic class `DeviceSplitKernel`. - Removes tile splitting, a new work pool implementation takes its place and allows as many threads as will fit in memory regardless of tile size, which can give performance gains. - Refactors split state buffers into one buffer, as well as reduces the number of arguments passed to kernels. Means there's less code to deal with overall. - Moves kernel logic out of OpenCL kernel files so they can later be used by other device types. - Replaced OpenCL specific APIs with new generic versions - Tiles can now be seen updating during rendering
2017-02-22 08:10:02 -05:00
RenderTile tile;
Cycles: avoid reallocating tile denoising memory many times during render.
2017-11-08 20:15:38 +01:00
DenoisingTask denoising(this);
Cycles: Split device_opencl.cpp into multiple files for easier maintenance There are no user-visible changes, just some internal restructuring. Differential Revision: https://developer.blender.org/D2231
2016-09-14 23:47:54 +02:00
Cycles: OpenCL split kernel refactor This does a few things at once: - Refactors host side split kernel logic into a new device agnostic class `DeviceSplitKernel`. - Removes tile splitting, a new work pool implementation takes its place and allows as many threads as will fit in memory regardless of tile size, which can give performance gains. - Refactors split state buffers into one buffer, as well as reduces the number of arguments passed to kernels. Means there's less code to deal with overall. - Moves kernel logic out of OpenCL kernel files so they can later be used by other device types. - Replaced OpenCL specific APIs with new generic versions - Tiles can now be seen updating during rendering
2017-02-22 08:10:02 -05:00
/* Allocate buffer for kernel globals */
Cycles: Fix compilation error without C++11
2017-10-24 11:14:01 +02:00
device_only_memory<KernelGlobalsDummy> kgbuffer(this, "kernel_globals");
Code refactor: move more memory allocation logic into device API. * Remove tex_* and pixels_* functions, replace by mem_*. * Add MEM_TEXTURE and MEM_PIXELS as memory types recognized by devices. * No longer create device_memory and call mem_* directly, always go through device_only_memory, device_vector and device_pixels.
2017-10-21 01:09:59 +02:00
kgbuffer.alloc_to_device(1);
Cycles: Split device_opencl.cpp into multiple files for easier maintenance There are no user-visible changes, just some internal restructuring. Differential Revision: https://developer.blender.org/D2231
2016-09-14 23:47:54 +02:00
Cycles: OpenCL split kernel refactor This does a few things at once: - Refactors host side split kernel logic into a new device agnostic class `DeviceSplitKernel`. - Removes tile splitting, a new work pool implementation takes its place and allows as many threads as will fit in memory regardless of tile size, which can give performance gains. - Refactors split state buffers into one buffer, as well as reduces the number of arguments passed to kernels. Means there's less code to deal with overall. - Moves kernel logic out of OpenCL kernel files so they can later be used by other device types. - Replaced OpenCL specific APIs with new generic versions - Tiles can now be seen updating during rendering
2017-02-22 08:10:02 -05:00
/* Keep rendering tiles until done. */
while(task->acquire_tile(this, tile)) {
Cycles: Implement denoising option for reducing noise in the rendered image This commit contains the first part of the new Cycles denoising option, which filters the resulting image using information gathered during rendering to get rid of noise while preserving visual features as well as possible. To use the option, enable it in the render layer options. The default settings fit a wide range of scenes, but the user can tweak individual settings to control the tradeoff between a noise-free image, image details, and calculation time. Note that the denoiser may still change in the future and that some features are not implemented yet. The most important missing feature is animation denoising, which uses information from multiple frames at once to produce a flicker-free and smoother result. These features will be added in the future. Finally, thanks to all the people who supported this project: - Google (through the GSoC) and Theory Studios for sponsoring the development - The authors of the papers I used for implementing the denoiser (more details on them will be included in the technical docs) - The other Cycles devs for feedback on the code, especially Sergey for mentoring the GSoC project and Brecht for the code review! - And of course the users who helped with testing, reported bugs and things that could and/or should work better!
2017-05-07 14:40:58 +02:00
if(tile.task == RenderTile::PATH_TRACE) {
assert(tile.task == RenderTile::PATH_TRACE);
split_kernel->path_trace(task,
tile,
kgbuffer,
*const_mem_map["__data"]);
/* 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;
Cycles: avoid reallocating tile denoising memory many times during render.
2017-11-08 20:15:38 +01:00
denoise(tile, denoising, *task);
Cycles: Implement denoising option for reducing noise in the rendered image This commit contains the first part of the new Cycles denoising option, which filters the resulting image using information gathered during rendering to get rid of noise while preserving visual features as well as possible. To use the option, enable it in the render layer options. The default settings fit a wide range of scenes, but the user can tweak individual settings to control the tradeoff between a noise-free image, image details, and calculation time. Note that the denoiser may still change in the future and that some features are not implemented yet. The most important missing feature is animation denoising, which uses information from multiple frames at once to produce a flicker-free and smoother result. These features will be added in the future. Finally, thanks to all the people who supported this project: - Google (through the GSoC) and Theory Studios for sponsoring the development - The authors of the papers I used for implementing the denoiser (more details on them will be included in the technical docs) - The other Cycles devs for feedback on the code, especially Sergey for mentoring the GSoC project and Brecht for the code review! - And of course the users who helped with testing, reported bugs and things that could and/or should work better!
2017-05-07 14:40:58 +02:00
task->update_progress(&tile, tile.w*tile.h);
}
Cycles: Split device_opencl.cpp into multiple files for easier maintenance There are no user-visible changes, just some internal restructuring. Differential Revision: https://developer.blender.org/D2231
2016-09-14 23:47:54 +02:00
Cycles: OpenCL split kernel refactor This does a few things at once: - Refactors host side split kernel logic into a new device agnostic class `DeviceSplitKernel`. - Removes tile splitting, a new work pool implementation takes its place and allows as many threads as will fit in memory regardless of tile size, which can give performance gains. - Refactors split state buffers into one buffer, as well as reduces the number of arguments passed to kernels. Means there's less code to deal with overall. - Moves kernel logic out of OpenCL kernel files so they can later be used by other device types. - Replaced OpenCL specific APIs with new generic versions - Tiles can now be seen updating during rendering
2017-02-22 08:10:02 -05:00
task->release_tile(tile);
}
Cycles: Split device_opencl.cpp into multiple files for easier maintenance There are no user-visible changes, just some internal restructuring. Differential Revision: https://developer.blender.org/D2231
2016-09-14 23:47:54 +02:00
Code refactor: move more memory allocation logic into device API. * Remove tex_* and pixels_* functions, replace by mem_*. * Add MEM_TEXTURE and MEM_PIXELS as memory types recognized by devices. * No longer create device_memory and call mem_* directly, always go through device_only_memory, device_vector and device_pixels.
2017-10-21 01:09:59 +02:00
kgbuffer.free();
Cycles: Split device_opencl.cpp into multiple files for easier maintenance There are no user-visible changes, just some internal restructuring. Differential Revision: https://developer.blender.org/D2231
2016-09-14 23:47:54 +02:00
}
}
Cycles: Implement denoising option for reducing noise in the rendered image This commit contains the first part of the new Cycles denoising option, which filters the resulting image using information gathered during rendering to get rid of noise while preserving visual features as well as possible. To use the option, enable it in the render layer options. The default settings fit a wide range of scenes, but the user can tweak individual settings to control the tradeoff between a noise-free image, image details, and calculation time. Note that the denoiser may still change in the future and that some features are not implemented yet. The most important missing feature is animation denoising, which uses information from multiple frames at once to produce a flicker-free and smoother result. These features will be added in the future. Finally, thanks to all the people who supported this project: - Google (through the GSoC) and Theory Studios for sponsoring the development - The authors of the papers I used for implementing the denoiser (more details on them will be included in the technical docs) - The other Cycles devs for feedback on the code, especially Sergey for mentoring the GSoC project and Brecht for the code review! - And of course the users who helped with testing, reported bugs and things that could and/or should work better!
2017-05-07 14:40:58 +02:00
bool is_split_kernel()
{
return true;
}
Cycles: OpenCL split kernel refactor This does a few things at once: - Refactors host side split kernel logic into a new device agnostic class `DeviceSplitKernel`. - Removes tile splitting, a new work pool implementation takes its place and allows as many threads as will fit in memory regardless of tile size, which can give performance gains. - Refactors split state buffers into one buffer, as well as reduces the number of arguments passed to kernels. Means there's less code to deal with overall. - Moves kernel logic out of OpenCL kernel files so they can later be used by other device types. - Replaced OpenCL specific APIs with new generic versions - Tiles can now be seen updating during rendering
2017-02-22 08:10:02 -05:00
protected:
/* ** Those guys are for workign around some compiler-specific bugs ** */
Cycles: Refactor Progress system to provide better estimates The Progress system in Cycles had two limitations so far: - It just counted tiles, but ignored their size. For example, when rendering a 600x500 image with 512x512 tiles, the right 88x500 tile would count for 50% of the progress, although it only covers 15% of the image. - Scene update time was incorrectly counted as rendering time - therefore, the remaining time started very long and gradually decreased. This patch fixes both problems: First of all, the Progress now has a function to ignore time spans, and that is used to ignore scene update time. The larger change is the tile size: Instead of counting samples per tile, so that the final value is num_samples*num_tiles, the code now counts every sample for every pixel, so that the final value is num_samples*num_pixels. Along with that, some unused variables were removed from the Progress and Session classes. Reviewers: brecht, sergey, #cycles Subscribers: brecht, candreacchio, sergey Differential Revision: https://developer.blender.org/D2214
2016-11-26 04:22:34 +01:00
Cycles: OpenCL split kernel refactor This does a few things at once: - Refactors host side split kernel logic into a new device agnostic class `DeviceSplitKernel`. - Removes tile splitting, a new work pool implementation takes its place and allows as many threads as will fit in memory regardless of tile size, which can give performance gains. - Refactors split state buffers into one buffer, as well as reduces the number of arguments passed to kernels. Means there's less code to deal with overall. - Moves kernel logic out of OpenCL kernel files so they can later be used by other device types. - Replaced OpenCL specific APIs with new generic versions - Tiles can now be seen updating during rendering
2017-02-22 08:10:02 -05:00
string build_options_for_base_program(
const DeviceRequestedFeatures& requested_features)
Cycles: Split device_opencl.cpp into multiple files for easier maintenance There are no user-visible changes, just some internal restructuring. Differential Revision: https://developer.blender.org/D2231
2016-09-14 23:47:54 +02:00
{
Cycles: OpenCL split kernel refactor This does a few things at once: - Refactors host side split kernel logic into a new device agnostic class `DeviceSplitKernel`. - Removes tile splitting, a new work pool implementation takes its place and allows as many threads as will fit in memory regardless of tile size, which can give performance gains. - Refactors split state buffers into one buffer, as well as reduces the number of arguments passed to kernels. Means there's less code to deal with overall. - Moves kernel logic out of OpenCL kernel files so they can later be used by other device types. - Replaced OpenCL specific APIs with new generic versions - Tiles can now be seen updating during rendering
2017-02-22 08:10:02 -05:00
return requested_features.get_build_options();
Cycles: Split device_opencl.cpp into multiple files for easier maintenance There are no user-visible changes, just some internal restructuring. Differential Revision: https://developer.blender.org/D2231
2016-09-14 23:47:54 +02:00
}
Cycles: OpenCL split kernel refactor This does a few things at once: - Refactors host side split kernel logic into a new device agnostic class `DeviceSplitKernel`. - Removes tile splitting, a new work pool implementation takes its place and allows as many threads as will fit in memory regardless of tile size, which can give performance gains. - Refactors split state buffers into one buffer, as well as reduces the number of arguments passed to kernels. Means there's less code to deal with overall. - Moves kernel logic out of OpenCL kernel files so they can later be used by other device types. - Replaced OpenCL specific APIs with new generic versions - Tiles can now be seen updating during rendering
2017-02-22 08:10:02 -05:00
friend class OpenCLSplitKernel;
friend class OpenCLSplitKernelFunction;
};
Cycles: Split device_opencl.cpp into multiple files for easier maintenance There are no user-visible changes, just some internal restructuring. Differential Revision: https://developer.blender.org/D2231
2016-09-14 23:47:54 +02:00
Cycles: Pass all buffers to each kernel call for OpenCL Technically not passing all buffers used by a kernel is undefined behavior. We haven't had any issues with this so far on AMD or Nvidia, but it's known to be a problem with Intel and we received a report from AMD that this is a problem on newer hardware, so we need to make this change at some point. Unfortunately there a cost to being correct, about 5% for the benchmark scenes. For low sample counts it's even worse, I've seen up to 50% slowdown. For the latter case I think adjusting tile updating logic can help, but not sure what that would look like yet (it would be just a few lines change however).
2017-05-11 19:23:49 -04:00
struct CachedSplitMemory {
int id;
device_memory *split_data;
device_memory *ray_state;
device_memory *queue_index;
device_memory *use_queues_flag;
device_memory *work_pools;
device_ptr *buffer;
};
Cycles: OpenCL split kernel refactor This does a few things at once: - Refactors host side split kernel logic into a new device agnostic class `DeviceSplitKernel`. - Removes tile splitting, a new work pool implementation takes its place and allows as many threads as will fit in memory regardless of tile size, which can give performance gains. - Refactors split state buffers into one buffer, as well as reduces the number of arguments passed to kernels. Means there's less code to deal with overall. - Moves kernel logic out of OpenCL kernel files so they can later be used by other device types. - Replaced OpenCL specific APIs with new generic versions - Tiles can now be seen updating during rendering
2017-02-22 08:10:02 -05:00
class OpenCLSplitKernelFunction : public SplitKernelFunction {
public:
OpenCLDeviceSplitKernel* device;
OpenCLDeviceBase::OpenCLProgram program;
Cycles: Pass all buffers to each kernel call for OpenCL Technically not passing all buffers used by a kernel is undefined behavior. We haven't had any issues with this so far on AMD or Nvidia, but it's known to be a problem with Intel and we received a report from AMD that this is a problem on newer hardware, so we need to make this change at some point. Unfortunately there a cost to being correct, about 5% for the benchmark scenes. For low sample counts it's even worse, I've seen up to 50% slowdown. For the latter case I think adjusting tile updating logic can help, but not sure what that would look like yet (it would be just a few lines change however).
2017-05-11 19:23:49 -04:00
CachedSplitMemory& cached_memory;
int cached_id;
OpenCLSplitKernelFunction(OpenCLDeviceSplitKernel* device, CachedSplitMemory& cached_memory) :
device(device), cached_memory(cached_memory), cached_id(cached_memory.id-1)
{
}
Cycles: Split device_opencl.cpp into multiple files for easier maintenance There are no user-visible changes, just some internal restructuring. Differential Revision: https://developer.blender.org/D2231
2016-09-14 23:47:54 +02:00
Cycles: Pass all buffers to each kernel call for OpenCL Technically not passing all buffers used by a kernel is undefined behavior. We haven't had any issues with this so far on AMD or Nvidia, but it's known to be a problem with Intel and we received a report from AMD that this is a problem on newer hardware, so we need to make this change at some point. Unfortunately there a cost to being correct, about 5% for the benchmark scenes. For low sample counts it's even worse, I've seen up to 50% slowdown. For the latter case I think adjusting tile updating logic can help, but not sure what that would look like yet (it would be just a few lines change however).
2017-05-11 19:23:49 -04:00
~OpenCLSplitKernelFunction()
{
program.release();
}
Cycles: Split device_opencl.cpp into multiple files for easier maintenance There are no user-visible changes, just some internal restructuring. Differential Revision: https://developer.blender.org/D2231
2016-09-14 23:47:54 +02:00
Cycles: OpenCL split kernel refactor This does a few things at once: - Refactors host side split kernel logic into a new device agnostic class `DeviceSplitKernel`. - Removes tile splitting, a new work pool implementation takes its place and allows as many threads as will fit in memory regardless of tile size, which can give performance gains. - Refactors split state buffers into one buffer, as well as reduces the number of arguments passed to kernels. Means there's less code to deal with overall. - Moves kernel logic out of OpenCL kernel files so they can later be used by other device types. - Replaced OpenCL specific APIs with new generic versions - Tiles can now be seen updating during rendering
2017-02-22 08:10:02 -05:00
virtual bool enqueue(const KernelDimensions& dim, device_memory& kg, device_memory& data)
Cycles: Split device_opencl.cpp into multiple files for easier maintenance There are no user-visible changes, just some internal restructuring. Differential Revision: https://developer.blender.org/D2231
2016-09-14 23:47:54 +02:00
{
Cycles: Pass all buffers to each kernel call for OpenCL Technically not passing all buffers used by a kernel is undefined behavior. We haven't had any issues with this so far on AMD or Nvidia, but it's known to be a problem with Intel and we received a report from AMD that this is a problem on newer hardware, so we need to make this change at some point. Unfortunately there a cost to being correct, about 5% for the benchmark scenes. For low sample counts it's even worse, I've seen up to 50% slowdown. For the latter case I think adjusting tile updating logic can help, but not sure what that would look like yet (it would be just a few lines change however).
2017-05-11 19:23:49 -04:00
if(cached_id != cached_memory.id) {
cl_uint start_arg_index =
device->kernel_set_args(program(),
0,
kg,
data,
*cached_memory.split_data,
Code refactor: remove rng_state buffer and compute hash on the fly. A little faster on some benchmark scenes, a little slower on others, seems about performance neutral on average and saves a little memory.
2017-09-27 00:39:53 +02:00
*cached_memory.ray_state);
Cycles: Pass all buffers to each kernel call for OpenCL Technically not passing all buffers used by a kernel is undefined behavior. We haven't had any issues with this so far on AMD or Nvidia, but it's known to be a problem with Intel and we received a report from AMD that this is a problem on newer hardware, so we need to make this change at some point. Unfortunately there a cost to being correct, about 5% for the benchmark scenes. For low sample counts it's even worse, I've seen up to 50% slowdown. For the latter case I think adjusting tile updating logic can help, but not sure what that would look like yet (it would be just a few lines change however).
2017-05-11 19:23:49 -04:00
Cycles: Pack kernel textures into buffers for OpenCL Image textures were being packed into a single buffer for OpenCL, which limited the amount of memory available for images to the size of one buffer (usually 4gb on AMD hardware). By packing textures into multiple buffers that limit is removed, while simultaneously reducing the number of buffers that need to be passed to each kernel. Benchmarks were within 2%. Fixes T51554. Differential Revision: https://developer.blender.org/D2745
2017-08-08 07:12:04 -04:00
device->set_kernel_arg_buffers(program(), &start_arg_index);
Cycles: Pass all buffers to each kernel call for OpenCL Technically not passing all buffers used by a kernel is undefined behavior. We haven't had any issues with this so far on AMD or Nvidia, but it's known to be a problem with Intel and we received a report from AMD that this is a problem on newer hardware, so we need to make this change at some point. Unfortunately there a cost to being correct, about 5% for the benchmark scenes. For low sample counts it's even worse, I've seen up to 50% slowdown. For the latter case I think adjusting tile updating logic can help, but not sure what that would look like yet (it would be just a few lines change however).
2017-05-11 19:23:49 -04:00
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;
}
Cycles: OpenCL split kernel refactor This does a few things at once: - Refactors host side split kernel logic into a new device agnostic class `DeviceSplitKernel`. - Removes tile splitting, a new work pool implementation takes its place and allows as many threads as will fit in memory regardless of tile size, which can give performance gains. - Refactors split state buffers into one buffer, as well as reduces the number of arguments passed to kernels. Means there's less code to deal with overall. - Moves kernel logic out of OpenCL kernel files so they can later be used by other device types. - Replaced OpenCL specific APIs with new generic versions - Tiles can now be seen updating during rendering
2017-02-22 08:10:02 -05:00
device->ciErr = clEnqueueNDRangeKernel(device->cqCommandQueue,
program(),
2,
NULL,
dim.global_size,
dim.local_size,
0,
NULL,
NULL);
device->opencl_assert_err(device->ciErr, "clEnqueueNDRangeKernel");
if(device->ciErr != CL_SUCCESS) {
string message = string_printf("OpenCL error: %s in clEnqueueNDRangeKernel()",
clewErrorString(device->ciErr));
device->opencl_error(message);
return false;
Cycles: Split device_opencl.cpp into multiple files for easier maintenance There are no user-visible changes, just some internal restructuring. Differential Revision: https://developer.blender.org/D2231
2016-09-14 23:47:54 +02:00
}
Cycles: OpenCL split kernel refactor This does a few things at once: - Refactors host side split kernel logic into a new device agnostic class `DeviceSplitKernel`. - Removes tile splitting, a new work pool implementation takes its place and allows as many threads as will fit in memory regardless of tile size, which can give performance gains. - Refactors split state buffers into one buffer, as well as reduces the number of arguments passed to kernels. Means there's less code to deal with overall. - Moves kernel logic out of OpenCL kernel files so they can later be used by other device types. - Replaced OpenCL specific APIs with new generic versions - Tiles can now be seen updating during rendering
2017-02-22 08:10:02 -05:00
return true;
}
};
Cycles: Split device_opencl.cpp into multiple files for easier maintenance There are no user-visible changes, just some internal restructuring. Differential Revision: https://developer.blender.org/D2231
2016-09-14 23:47:54 +02:00
Cycles: OpenCL split kernel refactor This does a few things at once: - Refactors host side split kernel logic into a new device agnostic class `DeviceSplitKernel`. - Removes tile splitting, a new work pool implementation takes its place and allows as many threads as will fit in memory regardless of tile size, which can give performance gains. - Refactors split state buffers into one buffer, as well as reduces the number of arguments passed to kernels. Means there's less code to deal with overall. - Moves kernel logic out of OpenCL kernel files so they can later be used by other device types. - Replaced OpenCL specific APIs with new generic versions - Tiles can now be seen updating during rendering
2017-02-22 08:10:02 -05:00
class OpenCLSplitKernel : public DeviceSplitKernel {
OpenCLDeviceSplitKernel *device;
Cycles: Pass all buffers to each kernel call for OpenCL Technically not passing all buffers used by a kernel is undefined behavior. We haven't had any issues with this so far on AMD or Nvidia, but it's known to be a problem with Intel and we received a report from AMD that this is a problem on newer hardware, so we need to make this change at some point. Unfortunately there a cost to being correct, about 5% for the benchmark scenes. For low sample counts it's even worse, I've seen up to 50% slowdown. For the latter case I think adjusting tile updating logic can help, but not sure what that would look like yet (it would be just a few lines change however).
2017-05-11 19:23:49 -04:00
CachedSplitMemory cached_memory;
Cycles: OpenCL split kernel refactor This does a few things at once: - Refactors host side split kernel logic into a new device agnostic class `DeviceSplitKernel`. - Removes tile splitting, a new work pool implementation takes its place and allows as many threads as will fit in memory regardless of tile size, which can give performance gains. - Refactors split state buffers into one buffer, as well as reduces the number of arguments passed to kernels. Means there's less code to deal with overall. - Moves kernel logic out of OpenCL kernel files so they can later be used by other device types. - Replaced OpenCL specific APIs with new generic versions - Tiles can now be seen updating during rendering
2017-02-22 08:10:02 -05:00
public:
explicit OpenCLSplitKernel(OpenCLDeviceSplitKernel *device) : DeviceSplitKernel(device), device(device) {
Cycles: Split device_opencl.cpp into multiple files for easier maintenance There are no user-visible changes, just some internal restructuring. Differential Revision: https://developer.blender.org/D2231
2016-09-14 23:47:54 +02:00
}
Cycles: Pass string by const reference rather than by value Some of the functions might have been inlined, but others i don't see how that was possible (don't think virtual functions can be inlined here). In any case, better be explicitly optimal in the code.
2017-07-05 12:27:41 +02:00
virtual SplitKernelFunction* get_split_kernel_function(const string& kernel_name,
Cycles: OpenCL split kernel refactor This does a few things at once: - Refactors host side split kernel logic into a new device agnostic class `DeviceSplitKernel`. - Removes tile splitting, a new work pool implementation takes its place and allows as many threads as will fit in memory regardless of tile size, which can give performance gains. - Refactors split state buffers into one buffer, as well as reduces the number of arguments passed to kernels. Means there's less code to deal with overall. - Moves kernel logic out of OpenCL kernel files so they can later be used by other device types. - Replaced OpenCL specific APIs with new generic versions - Tiles can now be seen updating during rendering
2017-02-22 08:10:02 -05:00
const DeviceRequestedFeatures& requested_features)
Cycles: Split device_opencl.cpp into multiple files for easier maintenance There are no user-visible changes, just some internal restructuring. Differential Revision: https://developer.blender.org/D2231
2016-09-14 23:47:54 +02:00
{
Cycles: Pass all buffers to each kernel call for OpenCL Technically not passing all buffers used by a kernel is undefined behavior. We haven't had any issues with this so far on AMD or Nvidia, but it's known to be a problem with Intel and we received a report from AMD that this is a problem on newer hardware, so we need to make this change at some point. Unfortunately there a cost to being correct, about 5% for the benchmark scenes. For low sample counts it's even worse, I've seen up to 50% slowdown. For the latter case I think adjusting tile updating logic can help, but not sure what that would look like yet (it would be just a few lines change however).
2017-05-11 19:23:49 -04:00
OpenCLSplitKernelFunction* kernel = new OpenCLSplitKernelFunction(device, cached_memory);
Cycles: OpenCL split kernel refactor This does a few things at once: - Refactors host side split kernel logic into a new device agnostic class `DeviceSplitKernel`. - Removes tile splitting, a new work pool implementation takes its place and allows as many threads as will fit in memory regardless of tile size, which can give performance gains. - Refactors split state buffers into one buffer, as well as reduces the number of arguments passed to kernels. Means there's less code to deal with overall. - Moves kernel logic out of OpenCL kernel files so they can later be used by other device types. - Replaced OpenCL specific APIs with new generic versions - Tiles can now be seen updating during rendering
2017-02-22 08:10:02 -05:00
Cycles: add single program debug option for split kernel Single program generally compiles kernels faster (2-3 times), loads faster, takes less drive space (2-3 times), and reduces the number of cached kernels.
2017-03-08 17:56:06 +01:00
bool single_program = OpenCLInfo::use_single_program();
kernel->program =
OpenCLDeviceBase::OpenCLProgram(device,
single_program ? "split" : "split_" + kernel_name,
single_program ? "kernel_split.cl" : "kernel_" + kernel_name + ".cl",
get_build_options(device, requested_features));
Cycles: OpenCL split kernel refactor This does a few things at once: - Refactors host side split kernel logic into a new device agnostic class `DeviceSplitKernel`. - Removes tile splitting, a new work pool implementation takes its place and allows as many threads as will fit in memory regardless of tile size, which can give performance gains. - Refactors split state buffers into one buffer, as well as reduces the number of arguments passed to kernels. Means there's less code to deal with overall. - Moves kernel logic out of OpenCL kernel files so they can later be used by other device types. - Replaced OpenCL specific APIs with new generic versions - Tiles can now be seen updating during rendering
2017-02-22 08:10:02 -05:00
kernel->program.add_kernel(ustring("path_trace_" + kernel_name));
kernel->program.load();
if(!kernel->program.is_loaded()) {
delete kernel;
return NULL;
Cycles: Split device_opencl.cpp into multiple files for easier maintenance There are no user-visible changes, just some internal restructuring. Differential Revision: https://developer.blender.org/D2231
2016-09-14 23:47:54 +02:00
}
Cycles: OpenCL split kernel refactor This does a few things at once: - Refactors host side split kernel logic into a new device agnostic class `DeviceSplitKernel`. - Removes tile splitting, a new work pool implementation takes its place and allows as many threads as will fit in memory regardless of tile size, which can give performance gains. - Refactors split state buffers into one buffer, as well as reduces the number of arguments passed to kernels. Means there's less code to deal with overall. - Moves kernel logic out of OpenCL kernel files so they can later be used by other device types. - Replaced OpenCL specific APIs with new generic versions - Tiles can now be seen updating during rendering
2017-02-22 08:10:02 -05:00
return kernel;
Cycles: Split device_opencl.cpp into multiple files for easier maintenance There are no user-visible changes, just some internal restructuring. Differential Revision: https://developer.blender.org/D2231
2016-09-14 23:47:54 +02:00
}
Fix T50888: Numeric overflow in split kernel state buffer size calculation Overflow led to the state buffer being too small and the split kernel to get stuck doing nothing forever.
2017-03-11 05:23:11 -05:00
virtual uint64_t state_buffer_size(device_memory& kg, device_memory& data, size_t num_threads)
Cycles: Calculate size of split state buffer kernel side By calculating the size of the state buffer in the kernel rather than the host less code is needed and the size actually reflects the requested features. Will also be a little faster in some cases because of larger global work size.
2017-03-04 06:29:01 -05:00
{
Code refactor: store device/interp/extension/type in each device_memory.
2017-10-20 23:31:13 +02:00
device_vector<uint64_t> size_buffer(device, "size_buffer", MEM_READ_WRITE);
Code refactor: move more memory allocation logic into device API. * Remove tex_* and pixels_* functions, replace by mem_*. * Add MEM_TEXTURE and MEM_PIXELS as memory types recognized by devices. * No longer create device_memory and call mem_* directly, always go through device_only_memory, device_vector and device_pixels.
2017-10-21 01:09:59 +02:00
size_buffer.alloc(1);
size_buffer.zero_to_device();
Cycles: Calculate size of split state buffer kernel side By calculating the size of the state buffer in the kernel rather than the host less code is needed and the size actually reflects the requested features. Will also be a little faster in some cases because of larger global work size.
2017-03-04 06:29:01 -05:00
uint threads = num_threads;
device->kernel_set_args(device->program_state_buffer_size(), 0, kg, data, threads, size_buffer);
size_t global_size = 64;
device->ciErr = clEnqueueNDRangeKernel(device->cqCommandQueue,
device->program_state_buffer_size(),
1,
NULL,
&global_size,
NULL,
0,
NULL,
NULL);
device->opencl_assert_err(device->ciErr, "clEnqueueNDRangeKernel");
Code refactor: move more memory allocation logic into device API. * Remove tex_* and pixels_* functions, replace by mem_*. * Add MEM_TEXTURE and MEM_PIXELS as memory types recognized by devices. * No longer create device_memory and call mem_* directly, always go through device_only_memory, device_vector and device_pixels.
2017-10-21 01:09:59 +02:00
size_buffer.copy_from_device(0, 1, 1);
Code refactor: avoid some unnecessary device memory copying.
2017-10-20 04:32:29 +02:00
size_t size = size_buffer[0];
Code refactor: move more memory allocation logic into device API. * Remove tex_* and pixels_* functions, replace by mem_*. * Add MEM_TEXTURE and MEM_PIXELS as memory types recognized by devices. * No longer create device_memory and call mem_* directly, always go through device_only_memory, device_vector and device_pixels.
2017-10-21 01:09:59 +02:00
size_buffer.free();
Cycles: Calculate size of split state buffer kernel side By calculating the size of the state buffer in the kernel rather than the host less code is needed and the size actually reflects the requested features. Will also be a little faster in some cases because of larger global work size.
2017-03-04 06:29:01 -05:00
if(device->ciErr != CL_SUCCESS) {
string message = string_printf("OpenCL error: %s in clEnqueueNDRangeKernel()",
clewErrorString(device->ciErr));
device->opencl_error(message);
return 0;
}
Code refactor: avoid some unnecessary device memory copying.
2017-10-20 04:32:29 +02:00
return size;
Cycles: Calculate size of split state buffer kernel side By calculating the size of the state buffer in the kernel rather than the host less code is needed and the size actually reflects the requested features. Will also be a little faster in some cases because of larger global work size.
2017-03-04 06:29:01 -05:00
}
Cycles: OpenCL split kernel refactor This does a few things at once: - Refactors host side split kernel logic into a new device agnostic class `DeviceSplitKernel`. - Removes tile splitting, a new work pool implementation takes its place and allows as many threads as will fit in memory regardless of tile size, which can give performance gains. - Refactors split state buffers into one buffer, as well as reduces the number of arguments passed to kernels. Means there's less code to deal with overall. - Moves kernel logic out of OpenCL kernel files so they can later be used by other device types. - Replaced OpenCL specific APIs with new generic versions - Tiles can now be seen updating during rendering
2017-02-22 08:10:02 -05:00
virtual bool enqueue_split_kernel_data_init(const KernelDimensions& dim,
RenderTile& rtile,
int num_global_elements,
device_memory& kernel_globals,
device_memory& kernel_data,
device_memory& split_data,
device_memory& ray_state,
device_memory& queue_index,
device_memory& use_queues_flag,
device_memory& work_pool_wgs
)
Cycles: Split device_opencl.cpp into multiple files for easier maintenance There are no user-visible changes, just some internal restructuring. Differential Revision: https://developer.blender.org/D2231
2016-09-14 23:47:54 +02:00
{
Cycles: OpenCL split kernel refactor This does a few things at once: - Refactors host side split kernel logic into a new device agnostic class `DeviceSplitKernel`. - Removes tile splitting, a new work pool implementation takes its place and allows as many threads as will fit in memory regardless of tile size, which can give performance gains. - Refactors split state buffers into one buffer, as well as reduces the number of arguments passed to kernels. Means there's less code to deal with overall. - Moves kernel logic out of OpenCL kernel files so they can later be used by other device types. - Replaced OpenCL specific APIs with new generic versions - Tiles can now be seen updating during rendering
2017-02-22 08:10:02 -05:00
cl_int dQueue_size = dim.global_size[0] * dim.global_size[1];
Cycles: Split device_opencl.cpp into multiple files for easier maintenance There are no user-visible changes, just some internal restructuring. Differential Revision: https://developer.blender.org/D2231
2016-09-14 23:47:54 +02:00
/* Set the range of samples to be processed for every ray in
* path-regeneration logic.
*/
cl_int start_sample = rtile.start_sample;
cl_int end_sample = rtile.start_sample + rtile.num_samples;
cl_uint start_arg_index =
Cycles: OpenCL split kernel refactor This does a few things at once: - Refactors host side split kernel logic into a new device agnostic class `DeviceSplitKernel`. - Removes tile splitting, a new work pool implementation takes its place and allows as many threads as will fit in memory regardless of tile size, which can give performance gains. - Refactors split state buffers into one buffer, as well as reduces the number of arguments passed to kernels. Means there's less code to deal with overall. - Moves kernel logic out of OpenCL kernel files so they can later be used by other device types. - Replaced OpenCL specific APIs with new generic versions - Tiles can now be seen updating during rendering
2017-02-22 08:10:02 -05:00
device->kernel_set_args(device->program_data_init(),
Cycles: Split device_opencl.cpp into multiple files for easier maintenance There are no user-visible changes, just some internal restructuring. Differential Revision: https://developer.blender.org/D2231
2016-09-14 23:47:54 +02:00
0,
Cycles: OpenCL split kernel refactor This does a few things at once: - Refactors host side split kernel logic into a new device agnostic class `DeviceSplitKernel`. - Removes tile splitting, a new work pool implementation takes its place and allows as many threads as will fit in memory regardless of tile size, which can give performance gains. - Refactors split state buffers into one buffer, as well as reduces the number of arguments passed to kernels. Means there's less code to deal with overall. - Moves kernel logic out of OpenCL kernel files so they can later be used by other device types. - Replaced OpenCL specific APIs with new generic versions - Tiles can now be seen updating during rendering
2017-02-22 08:10:02 -05:00
kernel_globals,
kernel_data,
Cycles: Fix indentation
2017-03-07 05:21:36 -05:00
split_data,
Cycles: OpenCL split kernel refactor This does a few things at once: - Refactors host side split kernel logic into a new device agnostic class `DeviceSplitKernel`. - Removes tile splitting, a new work pool implementation takes its place and allows as many threads as will fit in memory regardless of tile size, which can give performance gains. - Refactors split state buffers into one buffer, as well as reduces the number of arguments passed to kernels. Means there's less code to deal with overall. - Moves kernel logic out of OpenCL kernel files so they can later be used by other device types. - Replaced OpenCL specific APIs with new generic versions - Tiles can now be seen updating during rendering
2017-02-22 08:10:02 -05:00
num_global_elements,
Code refactor: remove rng_state buffer and compute hash on the fly. A little faster on some benchmark scenes, a little slower on others, seems about performance neutral on average and saves a little memory.
2017-09-27 00:39:53 +02:00
ray_state);
Cycles: Split device_opencl.cpp into multiple files for easier maintenance There are no user-visible changes, just some internal restructuring. Differential Revision: https://developer.blender.org/D2231
2016-09-14 23:47:54 +02:00
Cycles: Pack kernel textures into buffers for OpenCL Image textures were being packed into a single buffer for OpenCL, which limited the amount of memory available for images to the size of one buffer (usually 4gb on AMD hardware). By packing textures into multiple buffers that limit is removed, while simultaneously reducing the number of buffers that need to be passed to each kernel. Benchmarks were within 2%. Fixes T51554. Differential Revision: https://developer.blender.org/D2745
2017-08-08 07:12:04 -04:00
device->set_kernel_arg_buffers(device->program_data_init(), &start_arg_index);
Cycles: Split device_opencl.cpp into multiple files for easier maintenance There are no user-visible changes, just some internal restructuring. Differential Revision: https://developer.blender.org/D2231
2016-09-14 23:47:54 +02:00
start_arg_index +=
Cycles: OpenCL split kernel refactor This does a few things at once: - Refactors host side split kernel logic into a new device agnostic class `DeviceSplitKernel`. - Removes tile splitting, a new work pool implementation takes its place and allows as many threads as will fit in memory regardless of tile size, which can give performance gains. - Refactors split state buffers into one buffer, as well as reduces the number of arguments passed to kernels. Means there's less code to deal with overall. - Moves kernel logic out of OpenCL kernel files so they can later be used by other device types. - Replaced OpenCL specific APIs with new generic versions - Tiles can now be seen updating during rendering
2017-02-22 08:10:02 -05:00
device->kernel_set_args(device->program_data_init(),
Cycles: Split device_opencl.cpp into multiple files for easier maintenance There are no user-visible changes, just some internal restructuring. Differential Revision: https://developer.blender.org/D2231
2016-09-14 23:47:54 +02:00
start_arg_index,
start_sample,
Cycles: OpenCL split kernel refactor This does a few things at once: - Refactors host side split kernel logic into a new device agnostic class `DeviceSplitKernel`. - Removes tile splitting, a new work pool implementation takes its place and allows as many threads as will fit in memory regardless of tile size, which can give performance gains. - Refactors split state buffers into one buffer, as well as reduces the number of arguments passed to kernels. Means there's less code to deal with overall. - Moves kernel logic out of OpenCL kernel files so they can later be used by other device types. - Replaced OpenCL specific APIs with new generic versions - Tiles can now be seen updating during rendering
2017-02-22 08:10:02 -05:00
end_sample,
rtile.x,
rtile.y,
rtile.w,
rtile.h,
rtile.offset,
rtile.stride,
queue_index,
Cycles: Split device_opencl.cpp into multiple files for easier maintenance There are no user-visible changes, just some internal restructuring. Differential Revision: https://developer.blender.org/D2231
2016-09-14 23:47:54 +02:00
dQueue_size,
use_queues_flag,
work_pool_wgs,
Cycles: OpenCL split kernel refactor This does a few things at once: - Refactors host side split kernel logic into a new device agnostic class `DeviceSplitKernel`. - Removes tile splitting, a new work pool implementation takes its place and allows as many threads as will fit in memory regardless of tile size, which can give performance gains. - Refactors split state buffers into one buffer, as well as reduces the number of arguments passed to kernels. Means there's less code to deal with overall. - Moves kernel logic out of OpenCL kernel files so they can later be used by other device types. - Replaced OpenCL specific APIs with new generic versions - Tiles can now be seen updating during rendering
2017-02-22 08:10:02 -05:00
rtile.num_samples,
Cycles: Fix indentation
2017-03-07 05:21:36 -05:00
rtile.buffer);
Cycles: Split device_opencl.cpp into multiple files for easier maintenance There are no user-visible changes, just some internal restructuring. Differential Revision: https://developer.blender.org/D2231
2016-09-14 23:47:54 +02:00
/* Enqueue ckPathTraceKernel_data_init kernel. */
Cycles: OpenCL split kernel refactor This does a few things at once: - Refactors host side split kernel logic into a new device agnostic class `DeviceSplitKernel`. - Removes tile splitting, a new work pool implementation takes its place and allows as many threads as will fit in memory regardless of tile size, which can give performance gains. - Refactors split state buffers into one buffer, as well as reduces the number of arguments passed to kernels. Means there's less code to deal with overall. - Moves kernel logic out of OpenCL kernel files so they can later be used by other device types. - Replaced OpenCL specific APIs with new generic versions - Tiles can now be seen updating during rendering
2017-02-22 08:10:02 -05:00
device->ciErr = clEnqueueNDRangeKernel(device->cqCommandQueue,
device->program_data_init(),
2,
NULL,
dim.global_size,
dim.local_size,
0,
NULL,
NULL);
device->opencl_assert_err(device->ciErr, "clEnqueueNDRangeKernel");
if(device->ciErr != CL_SUCCESS) {
string message = string_printf("OpenCL error: %s in clEnqueueNDRangeKernel()",
clewErrorString(device->ciErr));
device->opencl_error(message);
Cycles: Split device_opencl.cpp into multiple files for easier maintenance There are no user-visible changes, just some internal restructuring. Differential Revision: https://developer.blender.org/D2231
2016-09-14 23:47:54 +02:00
return false;
}
Cycles: Pass all buffers to each kernel call for OpenCL Technically not passing all buffers used by a kernel is undefined behavior. We haven't had any issues with this so far on AMD or Nvidia, but it's known to be a problem with Intel and we received a report from AMD that this is a problem on newer hardware, so we need to make this change at some point. Unfortunately there a cost to being correct, about 5% for the benchmark scenes. For low sample counts it's even worse, I've seen up to 50% slowdown. For the latter case I think adjusting tile updating logic can help, but not sure what that would look like yet (it would be just a few lines change however).
2017-05-11 19:23:49 -04:00
cached_memory.split_data = &split_data;
cached_memory.ray_state = &ray_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++;
Cycles: OpenCL split kernel refactor This does a few things at once: - Refactors host side split kernel logic into a new device agnostic class `DeviceSplitKernel`. - Removes tile splitting, a new work pool implementation takes its place and allows as many threads as will fit in memory regardless of tile size, which can give performance gains. - Refactors split state buffers into one buffer, as well as reduces the number of arguments passed to kernels. Means there's less code to deal with overall. - Moves kernel logic out of OpenCL kernel files so they can later be used by other device types. - Replaced OpenCL specific APIs with new generic versions - Tiles can now be seen updating during rendering
2017-02-22 08:10:02 -05:00
return true;
Cycles: Split device_opencl.cpp into multiple files for easier maintenance There are no user-visible changes, just some internal restructuring. Differential Revision: https://developer.blender.org/D2231
2016-09-14 23:47:54 +02:00
}
Cycles: OpenCL split kernel refactor This does a few things at once: - Refactors host side split kernel logic into a new device agnostic class `DeviceSplitKernel`. - Removes tile splitting, a new work pool implementation takes its place and allows as many threads as will fit in memory regardless of tile size, which can give performance gains. - Refactors split state buffers into one buffer, as well as reduces the number of arguments passed to kernels. Means there's less code to deal with overall. - Moves kernel logic out of OpenCL kernel files so they can later be used by other device types. - Replaced OpenCL specific APIs with new generic versions - Tiles can now be seen updating during rendering
2017-02-22 08:10:02 -05:00
virtual int2 split_kernel_local_size()
Cycles: Split device_opencl.cpp into multiple files for easier maintenance There are no user-visible changes, just some internal restructuring. Differential Revision: https://developer.blender.org/D2231
2016-09-14 23:47:54 +02:00
{
Cycles: OpenCL split kernel refactor This does a few things at once: - Refactors host side split kernel logic into a new device agnostic class `DeviceSplitKernel`. - Removes tile splitting, a new work pool implementation takes its place and allows as many threads as will fit in memory regardless of tile size, which can give performance gains. - Refactors split state buffers into one buffer, as well as reduces the number of arguments passed to kernels. Means there's less code to deal with overall. - Moves kernel logic out of OpenCL kernel files so they can later be used by other device types. - Replaced OpenCL specific APIs with new generic versions - Tiles can now be seen updating during rendering
2017-02-22 08:10:02 -05:00
return make_int2(64, 1);
Cycles: Split device_opencl.cpp into multiple files for easier maintenance There are no user-visible changes, just some internal restructuring. Differential Revision: https://developer.blender.org/D2231
2016-09-14 23:47:54 +02:00
}
Cycles: Fix warning with MSVC
2017-04-07 18:28:16 +02:00
virtual int2 split_kernel_global_size(device_memory& kg, device_memory& data, DeviceTask * /*task*/)
Cycles: Split device_opencl.cpp into multiple files for easier maintenance There are no user-visible changes, just some internal restructuring. Differential Revision: https://developer.blender.org/D2231
2016-09-14 23:47:54 +02:00
{
Cycles: Simplify some extra OpenCL query code
2017-03-20 18:59:46 +01:00
cl_device_type type = OpenCLInfo::get_device_type(device->cdDevice);
Cycles: Adjust global size for OpenCL CPU devices to make them faster
2017-03-16 02:57:45 -04:00
/* Use small global size on CPU devices as it seems to be much faster. */
if(type == CL_DEVICE_TYPE_CPU) {
VLOG(1) << "Global size: (64, 64).";
return make_int2(64, 64);
}
Fix T50888: Numeric overflow in split kernel state buffer size calculation Overflow led to the state buffer being too small and the split kernel to get stuck doing nothing forever.
2017-03-11 05:23:11 -05:00
cl_ulong max_buffer_size;
clGetDeviceInfo(device->cdDevice, CL_DEVICE_MAX_MEM_ALLOC_SIZE, sizeof(cl_ulong), &max_buffer_size, NULL);
Cycles: Add artificial memory limit debug option for OpenCL
2017-07-05 20:16:41 -04:00
if(DebugFlags().opencl.mem_limit) {
Cycles: Fix ambiguity in call of min() function
2017-07-07 10:40:19 +02:00
max_buffer_size = min(max_buffer_size,
cl_ulong(DebugFlags().opencl.mem_limit - device->stats.mem_used));
Cycles: Add artificial memory limit debug option for OpenCL
2017-07-05 20:16:41 -04:00
}
Cycles: Adjust global size for OpenCL CPU devices to make them faster
2017-03-16 02:57:45 -04:00
VLOG(1) << "Maximum device allocation size: "
Cycles: Report device maximum allocation and detected global size
2017-02-20 13:02:11 +01:00
<< string_human_readable_number(max_buffer_size) << " bytes. ("
<< string_human_readable_size(max_buffer_size) << ").";
Cycles: Split device_opencl.cpp into multiple files for easier maintenance There are no user-visible changes, just some internal restructuring. Differential Revision: https://developer.blender.org/D2231
2016-09-14 23:47:54 +02:00
Cycles: Add another limit to OpenCL memory usage Some drivers may report very large allocation sizes, which could cause unnecessary memory usage. This is now limited to 2gb which should still be enough to get the needed performance benefits without waste.
2017-11-02 08:10:24 -04:00
/* Limit to 2gb, as we shouldn't need more than that and some devices may support much more. */
max_buffer_size = min(max_buffer_size / 2, (cl_ulong)2l*1024*1024*1024);
size_t num_elements = max_elements_for_max_buffer_size(kg, data, max_buffer_size);
Cycles: Dont allow global size to fall to zero
2017-07-05 20:15:47 -04:00
int2 global_size = make_int2(max(round_down((int)sqrt(num_elements), 64), 64), (int)sqrt(num_elements));
Cycles: Report device maximum allocation and detected global size
2017-02-20 13:02:11 +01:00
VLOG(1) << "Global size: " << global_size << ".";
Cycles: OpenCL split kernel refactor This does a few things at once: - Refactors host side split kernel logic into a new device agnostic class `DeviceSplitKernel`. - Removes tile splitting, a new work pool implementation takes its place and allows as many threads as will fit in memory regardless of tile size, which can give performance gains. - Refactors split state buffers into one buffer, as well as reduces the number of arguments passed to kernels. Means there's less code to deal with overall. - Moves kernel logic out of OpenCL kernel files so they can later be used by other device types. - Replaced OpenCL specific APIs with new generic versions - Tiles can now be seen updating during rendering
2017-02-22 08:10:02 -05:00
return global_size;
Cycles: Split device_opencl.cpp into multiple files for easier maintenance There are no user-visible changes, just some internal restructuring. Differential Revision: https://developer.blender.org/D2231
2016-09-14 23:47:54 +02:00
}
Cycles: OpenCL split kernel refactor This does a few things at once: - Refactors host side split kernel logic into a new device agnostic class `DeviceSplitKernel`. - Removes tile splitting, a new work pool implementation takes its place and allows as many threads as will fit in memory regardless of tile size, which can give performance gains. - Refactors split state buffers into one buffer, as well as reduces the number of arguments passed to kernels. Means there's less code to deal with overall. - Moves kernel logic out of OpenCL kernel files so they can later be used by other device types. - Replaced OpenCL specific APIs with new generic versions - Tiles can now be seen updating during rendering
2017-02-22 08:10:02 -05:00
};
Cycles: Split device_opencl.cpp into multiple files for easier maintenance There are no user-visible changes, just some internal restructuring. Differential Revision: https://developer.blender.org/D2231
2016-09-14 23:47:54 +02:00
Cycles: OpenCL split kernel refactor This does a few things at once: - Refactors host side split kernel logic into a new device agnostic class `DeviceSplitKernel`. - Removes tile splitting, a new work pool implementation takes its place and allows as many threads as will fit in memory regardless of tile size, which can give performance gains. - Refactors split state buffers into one buffer, as well as reduces the number of arguments passed to kernels. Means there's less code to deal with overall. - Moves kernel logic out of OpenCL kernel files so they can later be used by other device types. - Replaced OpenCL specific APIs with new generic versions - Tiles can now be seen updating during rendering
2017-02-22 08:10:02 -05:00
OpenCLDeviceSplitKernel::OpenCLDeviceSplitKernel(DeviceInfo& info, Stats &stats, bool background_)
: OpenCLDeviceBase(info, stats, background_)
{
split_kernel = new OpenCLSplitKernel(this);
Cycles: Split device_opencl.cpp into multiple files for easier maintenance There are no user-visible changes, just some internal restructuring. Differential Revision: https://developer.blender.org/D2231
2016-09-14 23:47:54 +02:00
Cycles: OpenCL split kernel refactor This does a few things at once: - Refactors host side split kernel logic into a new device agnostic class `DeviceSplitKernel`. - Removes tile splitting, a new work pool implementation takes its place and allows as many threads as will fit in memory regardless of tile size, which can give performance gains. - Refactors split state buffers into one buffer, as well as reduces the number of arguments passed to kernels. Means there's less code to deal with overall. - Moves kernel logic out of OpenCL kernel files so they can later be used by other device types. - Replaced OpenCL specific APIs with new generic versions - Tiles can now be seen updating during rendering
2017-02-22 08:10:02 -05:00
background = background_;
}
Cycles: Split device_opencl.cpp into multiple files for easier maintenance There are no user-visible changes, just some internal restructuring. Differential Revision: https://developer.blender.org/D2231
2016-09-14 23:47:54 +02:00
Device *opencl_create_split_device(DeviceInfo& info, Stats& stats, bool background)
{
return new OpenCLDeviceSplitKernel(info, stats, background);
}
CCL_NAMESPACE_END
#endif /* WITH_OPENCL */
Reference in New Issue Copy Permalink