mirror of
https://github.com/blender/blender
synced 2026-09-29 04:37:17 +03:00
f0ec6f183b
Both shadow kernels seems to be expensive enough and having them in the main module leads to very long module creation times. These new modules are always created (at least for now), but they are created in parallel with other modules. Overall it drastically reduces the time it takes to run OptiX OSL tests for the first time. Combined with the previous commits this reduces: - cycles_bake_optix_osl from 1330.44 sec to 177.30 sec - cycles_attributes_optix_osl from 2809.71 sec to 501.62 sec Measured on i9-11900K, NVIDIA RTX 6000 Ada. Pull Request: https://projects.blender.org/blender/blender/pulls/163706
89 lines
3.8 KiB
Text
89 lines
3.8 KiB
Text
/* SPDX-FileCopyrightText: 2011-2022 Blender Foundation
|
|
*
|
|
* SPDX-License-Identifier: Apache-2.0 */
|
|
|
|
#define WITH_OSL
|
|
|
|
/* Copy of the regular OptiX kernels with additional OSL support. */
|
|
|
|
#include "kernel/device/optix/kernel.cu"
|
|
|
|
#include "kernel/bake/bake.h"
|
|
#include "kernel/integrator/shade_background.h"
|
|
#include "kernel/integrator/shade_dedicated_light.h"
|
|
#include "kernel/integrator/shade_light.h"
|
|
#include "kernel/integrator/shade_surface.h"
|
|
|
|
#include "kernel/device/gpu/work_stealing.h"
|
|
|
|
extern "C" __global__ void __raygen__kernel_optix_integrator_shade_background()
|
|
{
|
|
const int global_index = optixGetLaunchIndex().x;
|
|
const int path_index = (kernel_params.path_index_array) ?
|
|
kernel_params.path_index_array[global_index] :
|
|
global_index;
|
|
integrator_shade_background(nullptr, path_index, kernel_params.render_buffer);
|
|
}
|
|
|
|
extern "C" __global__ void __raygen__kernel_optix_integrator_shade_light_nee()
|
|
{
|
|
const int global_index = optixGetLaunchIndex().x;
|
|
const int path_index = (kernel_params.path_index_array) ?
|
|
kernel_params.path_index_array[global_index] :
|
|
global_index;
|
|
integrator_shade_light_nee(nullptr, path_index, kernel_params.render_buffer);
|
|
}
|
|
|
|
extern "C" __global__ void __raygen__kernel_optix_integrator_shade_light_forward()
|
|
{
|
|
const int global_index = optixGetLaunchIndex().x;
|
|
const int path_index = (kernel_params.path_index_array) ?
|
|
kernel_params.path_index_array[global_index] :
|
|
global_index;
|
|
integrator_shade_light_forward(nullptr, path_index, kernel_params.render_buffer);
|
|
}
|
|
|
|
extern "C" __global__ void __raygen__kernel_optix_integrator_shade_surface()
|
|
{
|
|
const int global_index = optixGetLaunchIndex().x;
|
|
const int path_index = (kernel_params.path_index_array) ?
|
|
kernel_params.path_index_array[global_index] :
|
|
global_index;
|
|
integrator_shade_surface(nullptr, path_index, kernel_params.render_buffer);
|
|
}
|
|
|
|
extern "C" __global__ void __raygen__kernel_optix_integrator_shade_dedicated_light()
|
|
{
|
|
const int global_index = optixGetLaunchIndex().x;
|
|
const int path_index = (kernel_params.path_index_array) ?
|
|
kernel_params.path_index_array[global_index] :
|
|
global_index;
|
|
integrator_shade_dedicated_light(nullptr, path_index, kernel_params.render_buffer);
|
|
}
|
|
|
|
extern "C" __global__ void __raygen__kernel_optix_shader_eval_displace()
|
|
{
|
|
KernelShaderEvalInput *const input = (KernelShaderEvalInput *)kernel_params.path_index_array;
|
|
float *const output = kernel_params.render_buffer;
|
|
uint *const cache_miss = kernel_params.shader_eval_cache_miss;
|
|
const int global_index = kernel_params.shader_eval_offset + optixGetLaunchIndex().x;
|
|
kernel_displace_evaluate(nullptr, input, output, cache_miss, global_index);
|
|
}
|
|
|
|
extern "C" __global__ void __raygen__kernel_optix_shader_eval_background()
|
|
{
|
|
KernelShaderEvalInput *const input = (KernelShaderEvalInput *)kernel_params.path_index_array;
|
|
float *const output = kernel_params.render_buffer;
|
|
uint *const cache_miss = kernel_params.shader_eval_cache_miss;
|
|
const int global_index = kernel_params.shader_eval_offset + optixGetLaunchIndex().x;
|
|
kernel_background_evaluate(nullptr, input, output, cache_miss, global_index);
|
|
}
|
|
|
|
extern "C" __global__ void __raygen__kernel_optix_shader_eval_volume_density()
|
|
{
|
|
KernelShaderEvalInput *const input = (KernelShaderEvalInput *)kernel_params.path_index_array;
|
|
float *const output = kernel_params.render_buffer;
|
|
uint *const cache_miss = kernel_params.shader_eval_cache_miss;
|
|
const int global_index = kernel_params.shader_eval_offset + optixGetLaunchIndex().x;
|
|
kernel_volume_density_evaluate(nullptr, input, output, cache_miss, global_index);
|
|
}
|