diff options
| author | Patrick Mours <pmours@nvidia.com> | 2022-11-09 16:25:32 +0300 |
|---|---|---|
| committer | Patrick Mours <pmours@nvidia.com> | 2022-11-09 17:30:21 +0300 |
| commit | e6b38deb9dbb58118f6ee644409ce52f06eac5e5 (patch) | |
| tree | aa1c384db146094482f24c94f704742c6624db00 | |
| parent | efe073f57c34b438d21750795e97458a3d007be7 (diff) | |
Cycles: Add basic support for using OSL with OptiX
This patch generalizes the OSL support in Cycles to include GPU
device types and adds an implementation for that in the OptiX
device. There are some caveats still, including simplified texturing
due to lack of OIIO on the GPU and a few missing OSL intrinsics.
Note that this is incomplete and missing an update to the OSL
library before being enabled! The implementation is already
committed now to simplify further development.
Maniphest Tasks: T101222
Differential Revision: https://developer.blender.org/D15902
43 files changed, 3478 insertions, 564 deletions
diff --git a/build_files/cmake/platform/platform_win32.cmake b/build_files/cmake/platform/platform_win32.cmake index 7a2d3ad948a..47673794652 100644 --- a/build_files/cmake/platform/platform_win32.cmake +++ b/build_files/cmake/platform/platform_win32.cmake @@ -419,7 +419,7 @@ if(WITH_IMAGE_OPENEXR) warn_hardcoded_paths(OpenEXR) set(OPENEXR ${LIBDIR}/openexr) set(OPENEXR_INCLUDE_DIR ${OPENEXR}/include) - set(OPENEXR_INCLUDE_DIRS ${OPENEXR_INCLUDE_DIR} ${IMATH_INCLUDE_DIRS} ${OPENEXR}/include/OpenEXR) + set(OPENEXR_INCLUDE_DIRS ${OPENEXR_INCLUDE_DIR} ${IMATH_INCLUDE_DIRS} ${OPENEXR_INCLUDE_DIR}/OpenEXR) set(OPENEXR_LIBPATH ${OPENEXR}/lib) # Check if the 3.x library name exists # if not assume this is a 2.x library folder @@ -568,7 +568,8 @@ if(WITH_OPENIMAGEIO) if(NOT OpenImageIO_FOUND) set(OPENIMAGEIO ${LIBDIR}/OpenImageIO) set(OPENIMAGEIO_LIBPATH ${OPENIMAGEIO}/lib) - set(OPENIMAGEIO_INCLUDE_DIRS ${OPENIMAGEIO}/include) + set(OPENIMAGEIO_INCLUDE_DIR ${OPENIMAGEIO}/include) + set(OPENIMAGEIO_INCLUDE_DIRS ${OPENIMAGEIO_INCLUDE_DIR}) set(OIIO_OPTIMIZED optimized ${OPENIMAGEIO_LIBPATH}/OpenImageIO.lib optimized ${OPENIMAGEIO_LIBPATH}/OpenImageIO_Util.lib) set(OIIO_DEBUG debug ${OPENIMAGEIO_LIBPATH}/OpenImageIO_d.lib debug ${OPENIMAGEIO_LIBPATH}/OpenImageIO_Util_d.lib) set(OPENIMAGEIO_LIBRARIES ${OIIO_OPTIMIZED} ${OIIO_DEBUG}) @@ -785,6 +786,14 @@ if(WITH_CYCLES AND WITH_CYCLES_OSL) endif() find_path(OSL_INCLUDE_DIR OSL/oslclosure.h PATHS ${CYCLES_OSL}/include) find_program(OSL_COMPILER NAMES oslc PATHS ${CYCLES_OSL}/bin) + file(STRINGS "${OSL_INCLUDE_DIR}/OSL/oslversion.h" OSL_LIBRARY_VERSION_MAJOR + REGEX "^[ \t]*#define[ \t]+OSL_LIBRARY_VERSION_MAJOR[ \t]+[0-9]+.*$") + file(STRINGS "${OSL_INCLUDE_DIR}/OSL/oslversion.h" OSL_LIBRARY_VERSION_MINOR + REGEX "^[ \t]*#define[ \t]+OSL_LIBRARY_VERSION_MINOR[ \t]+[0-9]+.*$") + string(REGEX REPLACE ".*#define[ \t]+OSL_LIBRARY_VERSION_MAJOR[ \t]+([.0-9]+).*" + "\\1" OSL_LIBRARY_VERSION_MAJOR ${OSL_LIBRARY_VERSION_MAJOR}) + string(REGEX REPLACE ".*#define[ \t]+OSL_LIBRARY_VERSION_MINOR[ \t]+([.0-9]+).*" + "\\1" OSL_LIBRARY_VERSION_MINOR ${OSL_LIBRARY_VERSION_MINOR}) endif() if(WITH_CYCLES AND WITH_CYCLES_EMBREE) diff --git a/intern/cycles/blender/addon/__init__.py b/intern/cycles/blender/addon/__init__.py index 05f27bdbd4d..354c9c23a53 100644 --- a/intern/cycles/blender/addon/__init__.py +++ b/intern/cycles/blender/addon/__init__.py @@ -58,7 +58,7 @@ class CyclesRender(bpy.types.RenderEngine): if not self.session: if self.is_preview: cscene = bpy.context.scene.cycles - use_osl = cscene.shading_system and cscene.device == 'CPU' + use_osl = cscene.shading_system engine.create(self, data, preview_osl=use_osl) else: diff --git a/intern/cycles/blender/addon/engine.py b/intern/cycles/blender/addon/engine.py index e33891fa7a2..83dc6332f47 100644 --- a/intern/cycles/blender/addon/engine.py +++ b/intern/cycles/blender/addon/engine.py @@ -155,6 +155,10 @@ def with_osl(): import _cycles return _cycles.with_osl +def osl_version(): + import _cycles + return _cycles.osl_version + def with_path_guiding(): import _cycles diff --git a/intern/cycles/blender/addon/properties.py b/intern/cycles/blender/addon/properties.py index f5cd88f6b6a..9d7c71417f2 100644 --- a/intern/cycles/blender/addon/properties.py +++ b/intern/cycles/blender/addon/properties.py @@ -290,7 +290,7 @@ class CyclesRenderSettings(bpy.types.PropertyGroup): ) shading_system: BoolProperty( name="Open Shading Language", - description="Use Open Shading Language (CPU rendering only)", + description="Use Open Shading Language", ) preview_pause: BoolProperty( diff --git a/intern/cycles/blender/addon/ui.py b/intern/cycles/blender/addon/ui.py index 305accc8f1a..11fa2bc62fb 100644 --- a/intern/cycles/blender/addon/ui.py +++ b/intern/cycles/blender/addon/ui.py @@ -2305,7 +2305,7 @@ def draw_device(self, context): col.prop(cscene, "device") from . import engine - if engine.with_osl() and use_cpu(context): + if engine.with_osl() and (use_cpu(context) or (use_optix(context) and (engine.osl_version()[1] >= 13 or engine.osl_version()[0] > 1))): col.prop(cscene, "shading_system") diff --git a/intern/cycles/device/device.h b/intern/cycles/device/device.h index 2e4d18241cf..06a2f5c7b01 100644 --- a/intern/cycles/device/device.h +++ b/intern/cycles/device/device.h @@ -160,6 +160,11 @@ class Device { return true; } + virtual bool load_osl_kernels() + { + return true; + } + /* GPU device only functions. * These may not be used on CPU or multi-devices. */ diff --git a/intern/cycles/device/kernel.cpp b/intern/cycles/device/kernel.cpp index 96a99cd62cd..27ca0d81817 100644 --- a/intern/cycles/device/kernel.cpp +++ b/intern/cycles/device/kernel.cpp @@ -7,6 +7,30 @@ CCL_NAMESPACE_BEGIN +bool device_kernel_has_shading(DeviceKernel kernel) +{ + return (kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_BACKGROUND || + kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT || + kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE || + kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE || + kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_MNEE || + kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_VOLUME || + kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_SHADOW || + kernel == DEVICE_KERNEL_SHADER_EVAL_DISPLACE || + kernel == DEVICE_KERNEL_SHADER_EVAL_BACKGROUND || + kernel == DEVICE_KERNEL_SHADER_EVAL_CURVE_SHADOW_TRANSPARENCY); +} + +bool device_kernel_has_intersection(DeviceKernel kernel) +{ + return (kernel == DEVICE_KERNEL_INTEGRATOR_INTERSECT_CLOSEST || + kernel == DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW || + kernel == DEVICE_KERNEL_INTEGRATOR_INTERSECT_SUBSURFACE || + kernel == DEVICE_KERNEL_INTEGRATOR_INTERSECT_VOLUME_STACK || + kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE || + kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_MNEE); +} + const char *device_kernel_as_string(DeviceKernel kernel) { switch (kernel) { diff --git a/intern/cycles/device/kernel.h b/intern/cycles/device/kernel.h index 4ae461f1f67..b829a891260 100644 --- a/intern/cycles/device/kernel.h +++ b/intern/cycles/device/kernel.h @@ -11,6 +11,9 @@ CCL_NAMESPACE_BEGIN +bool device_kernel_has_shading(DeviceKernel kernel); +bool device_kernel_has_intersection(DeviceKernel kernel); + const char *device_kernel_as_string(DeviceKernel kernel); std::ostream &operator<<(std::ostream &os, DeviceKernel kernel); diff --git a/intern/cycles/device/multi/device.cpp b/intern/cycles/device/multi/device.cpp index 6904d2c2dc6..9605c6a7538 100644 --- a/intern/cycles/device/multi/device.cpp +++ b/intern/cycles/device/multi/device.cpp @@ -138,6 +138,15 @@ class MultiDevice : public Device { return true; } + bool load_osl_kernels() override + { + foreach (SubDevice &sub, devices) + if (!sub.device->load_osl_kernels()) + return false; + + return true; + } + void build_bvh(BVH *bvh, Progress &progress, bool refit) override { /* Try to build and share a single acceleration structure, if possible */ @@ -204,10 +213,12 @@ class MultiDevice : public Device { virtual void *get_cpu_osl_memory() override { - if (devices.size() > 1) { + /* Always return the OSL memory of the CPU device (this works since the constructor above + * guarantees that CPU devices are always added to the back). */ + if (devices.size() > 1 && devices.back().device->info.type != DEVICE_CPU) { return NULL; } - return devices.front().device->get_cpu_osl_memory(); + return devices.back().device->get_cpu_osl_memory(); } bool is_resident(device_ptr key, Device *sub_device) override diff --git a/intern/cycles/device/optix/device.cpp b/intern/cycles/device/optix/device.cpp index 68ca21374fd..58b72374a7d 100644 --- a/intern/cycles/device/optix/device.cpp +++ b/intern/cycles/device/optix/device.cpp @@ -9,6 +9,10 @@ #include "util/log.h" +#ifdef WITH_OSL +# include <OSL/oslversion.h> +#endif + #ifdef WITH_OPTIX # include <optix_function_table_definition.h> #endif @@ -65,6 +69,9 @@ void device_optix_info(const vector<DeviceInfo> &cuda_devices, vector<DeviceInfo info.type = DEVICE_OPTIX; info.id += "_OptiX"; +# if defined(WITH_OSL) && (OSL_VERSION_MINOR >= 13 || OSL_VERSION_MAJOR > 1) + info.has_osl = true; +# endif info.denoisers |= DENOISER_OPTIX; devices.push_back(info); diff --git a/intern/cycles/device/optix/device_impl.cpp b/intern/cycles/device/optix/device_impl.cpp index fabf4d7b69d..02f34bf3bd0 100644 --- a/intern/cycles/device/optix/device_impl.cpp +++ b/intern/cycles/device/optix/device_impl.cpp @@ -312,16 +312,34 @@ OptiXDevice::~OptiXDevice() if (optix_module != NULL) { optixModuleDestroy(optix_module); } - for (unsigned int i = 0; i < 2; ++i) { + for (int i = 0; i < 2; ++i) { if (builtin_modules[i] != NULL) { optixModuleDestroy(builtin_modules[i]); } } - for (unsigned int i = 0; i < NUM_PIPELINES; ++i) { + for (int i = 0; i < NUM_PIPELINES; ++i) { if (pipelines[i] != NULL) { optixPipelineDestroy(pipelines[i]); } } + for (int i = 0; i < NUM_PROGRAM_GROUPS; ++i) { + if (groups[i] != NULL) { + optixProgramGroupDestroy(groups[i]); + } + } + +# ifdef WITH_OSL + for (const OptixModule &module : osl_modules) { + if (module != NULL) { + optixModuleDestroy(module); + } + } + for (const OptixProgramGroup &group : osl_groups) { + if (group != NULL) { + optixProgramGroupDestroy(group); + } + } +# endif /* Make sure denoiser is destroyed before device context! */ if (denoiser_.optix_denoiser != nullptr) { @@ -381,6 +399,12 @@ bool OptiXDevice::load_kernels(const uint kernel_features) return false; } +# ifdef WITH_OSL + const bool use_osl = (kernel_features & KERNEL_FEATURE_OSL); +# else + const bool use_osl = false; +# endif + /* Skip creating OptiX module if only doing denoising. */ const bool need_optix_kernels = (kernel_features & (KERNEL_FEATURE_PATH_TRACING | KERNEL_FEATURE_BAKING)); @@ -388,12 +412,13 @@ bool OptiXDevice::load_kernels(const uint kernel_features) /* Detect existence of OptiX kernel and SDK here early. So we can error out * before compiling the CUDA kernels, to avoid failing right after when * compiling the OptiX kernel. */ + string suffix = use_osl ? "_osl" : + (kernel_features & (KERNEL_FEATURE_NODE_RAYTRACE | KERNEL_FEATURE_MNEE)) ? + "_shader_raytrace" : + ""; string ptx_filename; if (need_optix_kernels) { - ptx_filename = path_get( - (kernel_features & (KERNEL_FEATURE_NODE_RAYTRACE | KERNEL_FEATURE_MNEE)) ? - "lib/kernel_optix_shader_raytrace.ptx" : - "lib/kernel_optix.ptx"); + ptx_filename = path_get("lib/kernel_optix" + suffix + ".ptx"); if (use_adaptive_compilation() || path_file_size(ptx_filename) == -1) { std::string optix_include_dir = get_optix_include_dir(); if (optix_include_dir.empty()) { @@ -429,18 +454,41 @@ bool OptiXDevice::load_kernels(const uint kernel_features) optixModuleDestroy(optix_module); optix_module = NULL; } - for (unsigned int i = 0; i < 2; ++i) { + for (int i = 0; i < 2; ++i) { if (builtin_modules[i] != NULL) { optixModuleDestroy(builtin_modules[i]); builtin_modules[i] = NULL; } } - for (unsigned int i = 0; i < NUM_PIPELINES; ++i) { + for (int i = 0; i < NUM_PIPELINES; ++i) { if (pipelines[i] != NULL) { optixPipelineDestroy(pipelines[i]); pipelines[i] = NULL; } } + for (int i = 0; i < NUM_PROGRAM_GROUPS; ++i) { + if (groups[i] != NULL) { + optixProgramGroupDestroy(groups[i]); + groups[i] = NULL; + } + } + +# ifdef WITH_OSL + /* Recreating base OptiX module invalidates all OSL modules too, since they link against it. */ + for (const OptixModule &module : osl_modules) { + if (module != NULL) { + optixModuleDestroy(module); + } + } + osl_modules.clear(); + + for (const OptixProgramGroup &group : osl_groups) { + if (group != NULL) { + optixProgramGroupDestroy(group); + } + } + osl_groups.clear(); +# endif OptixModuleCompileOptions module_options = {}; module_options.maxRegisterCount = 0; /* Do not set an explicit register limit. */ @@ -461,7 +509,6 @@ bool OptiXDevice::load_kernels(const uint kernel_features) module_options.numPayloadTypes = 0; # endif - OptixPipelineCompileOptions pipeline_options = {}; /* Default to no motion blur and two-level graph, since it is the fastest option. */ pipeline_options.usesMotionBlur = false; pipeline_options.traversableGraphFlags = @@ -490,9 +537,7 @@ bool OptiXDevice::load_kernels(const uint kernel_features) /* Keep track of whether motion blur is enabled, so to enable/disable motion in BVH builds * This is necessary since objects may be reported to have motion if the Vector pass is * active, but may still need to be rendered without motion blur if that isn't active as well. */ - motion_blur = (kernel_features & KERNEL_FEATURE_OBJECT_MOTION) != 0; - - if (motion_blur) { + if (kernel_features & KERNEL_FEATURE_OBJECT_MOTION) { pipeline_options.usesMotionBlur = true; /* Motion blur can insert motion transforms into the traversal graph. * It is no longer a two-level graph then, so need to set flags to allow any configuration. */ @@ -503,13 +548,7 @@ bool OptiXDevice::load_kernels(const uint kernel_features) string ptx_data; if (use_adaptive_compilation() || path_file_size(ptx_filename) == -1) { string cflags = compile_kernel_get_common_cflags(kernel_features); - ptx_filename = compile_kernel( - cflags, - (kernel_features & (KERNEL_FEATURE_NODE_RAYTRACE | KERNEL_FEATURE_MNEE)) ? - "kernel_shader_raytrace" : - "kernel", - "optix", - true); + ptx_filename = compile_kernel(cflags, ("kernel" + suffix).c_str(), "optix", true); } if (ptx_filename.empty() || !path_read_text(ptx_filename, ptx_data)) { set_error(string_printf("Failed to load OptiX kernel from '%s'", ptx_filename.c_str())); @@ -551,7 +590,6 @@ bool OptiXDevice::load_kernels(const uint kernel_features) } /* Create program groups. */ - OptixProgramGroup groups[NUM_PROGRAM_GROUPS] = {}; OptixProgramGroupDesc group_descs[NUM_PROGRAM_GROUPS] = {}; OptixProgramGroupOptions group_options = {}; /* There are no options currently. */ group_descs[PG_RGEN_INTERSECT_CLOSEST].kind = OPTIX_PROGRAM_GROUP_KIND_RAYGEN; @@ -609,7 +647,7 @@ bool OptiXDevice::load_kernels(const uint kernel_features) group_descs[PG_HITS].hitgroup.moduleIS = builtin_modules[0]; group_descs[PG_HITS].hitgroup.entryFunctionNameIS = nullptr; - if (motion_blur) { + if (pipeline_options.usesMotionBlur) { builtin_options.usesMotionBlur = true; optix_assert(optixBuiltinISModuleGet( @@ -630,7 +668,6 @@ bool OptiXDevice::load_kernels(const uint kernel_features) } } - /* Pointclouds */ if (kernel_features & KERNEL_FEATURE_POINTCLOUD) { group_descs[PG_HITD_POINTCLOUD] = group_descs[PG_HITD]; group_descs[PG_HITD_POINTCLOUD].kind = OPTIX_PROGRAM_GROUP_KIND_HITGROUP; @@ -642,8 +679,8 @@ bool OptiXDevice::load_kernels(const uint kernel_features) group_descs[PG_HITS_POINTCLOUD].hitgroup.entryFunctionNameIS = "__intersection__point"; } + /* Add hit group for local intersections. */ if (kernel_features & (KERNEL_FEATURE_SUBSURFACE | KERNEL_FEATURE_NODE_RAYTRACE)) { - /* Add hit group for local intersections. */ group_descs[PG_HITL].kind = OPTIX_PROGRAM_GROUP_KIND_HITGROUP; group_descs[PG_HITL].hitgroup.moduleAH = optix_module; group_descs[PG_HITL].hitgroup.entryFunctionNameAH = "__anyhit__kernel_optix_local_hit"; @@ -655,16 +692,19 @@ bool OptiXDevice::load_kernels(const uint kernel_features) group_descs[PG_RGEN_SHADE_SURFACE_RAYTRACE].raygen.module = optix_module; group_descs[PG_RGEN_SHADE_SURFACE_RAYTRACE].raygen.entryFunctionName = "__raygen__kernel_optix_integrator_shade_surface_raytrace"; - group_descs[PG_CALL_SVM_AO].kind = OPTIX_PROGRAM_GROUP_KIND_CALLABLES; - group_descs[PG_CALL_SVM_AO].callables.moduleDC = optix_module; - group_descs[PG_CALL_SVM_AO].callables.entryFunctionNameDC = "__direct_callable__svm_node_ao"; - group_descs[PG_CALL_SVM_BEVEL].kind = OPTIX_PROGRAM_GROUP_KIND_CALLABLES; - group_descs[PG_CALL_SVM_BEVEL].callables.moduleDC = optix_module; - group_descs[PG_CALL_SVM_BEVEL].callables.entryFunctionNameDC = - "__direct_callable__svm_node_bevel"; + + /* Kernels with OSL support are built without SVM, so can skip those direct callables there. */ + if (!use_osl) { + group_descs[PG_CALL_SVM_AO].kind = OPTIX_PROGRAM_GROUP_KIND_CALLABLES; + group_descs[PG_CALL_SVM_AO].callables.moduleDC = optix_module; + group_descs[PG_CALL_SVM_AO].callables.entryFunctionNameDC = "__direct_callable__svm_node_ao"; + group_descs[PG_CALL_SVM_BEVEL].kind = OPTIX_PROGRAM_GROUP_KIND_CALLABLES; + group_descs[PG_CALL_SVM_BEVEL].callables.moduleDC = optix_module; + group_descs[PG_CALL_SVM_BEVEL].callables.entryFunctionNameDC = + "__direct_callable__svm_node_bevel"; + } } - /* MNEE. */ if (kernel_features & KERNEL_FEATURE_MNEE) { group_descs[PG_RGEN_SHADE_SURFACE_MNEE].kind = OPTIX_PROGRAM_GROUP_KIND_RAYGEN; group_descs[PG_RGEN_SHADE_SURFACE_MNEE].raygen.module = optix_module; @@ -672,6 +712,42 @@ bool OptiXDevice::load_kernels(const uint kernel_features) "__raygen__kernel_optix_integrator_shade_surface_mnee"; } + /* OSL uses direct callables to execute, so shading needs to be done in OptiX if OSL is used. */ + if (use_osl) { + group_descs[PG_RGEN_SHADE_BACKGROUND].kind = OPTIX_PROGRAM_GROUP_KIND_RAYGEN; + group_descs[PG_RGEN_SHADE_BACKGROUND].raygen.module = optix_module; + group_descs[PG_RGEN_SHADE_BACKGROUND].raygen.entryFunctionName = + "__raygen__kernel_optix_integrator_shade_background"; + group_descs[PG_RGEN_SHADE_LIGHT].kind = OPTIX_PROGRAM_GROUP_KIND_RAYGEN; + group_descs[PG_RGEN_SHADE_LIGHT].raygen.module = optix_module; + group_descs[PG_RGEN_SHADE_LIGHT].raygen.entryFunctionName = + "__raygen__kernel_optix_integrator_shade_light"; + group_descs[PG_RGEN_SHADE_SURFACE].kind = OPTIX_PROGRAM_GROUP_KIND_RAYGEN; + group_descs[PG_RGEN_SHADE_SURFACE].raygen.module = optix_module; + group_descs[PG_RGEN_SHADE_SURFACE].raygen.entryFunctionName = + "__raygen__kernel_optix_integrator_shade_surface"; + group_descs[PG_RGEN_SHADE_VOLUME].kind = OPTIX_PROGRAM_GROUP_KIND_RAYGEN; + group_descs[PG_RGEN_SHADE_VOLUME].raygen.module = optix_module; + group_descs[PG_RGEN_SHADE_VOLUME].raygen.entryFunctionName = + "__raygen__kernel_optix_integrator_shade_volume"; + group_descs[PG_RGEN_SHADE_SHADOW].kind = OPTIX_PROGRAM_GROUP_KIND_RAYGEN; + group_descs[PG_RGEN_SHADE_SHADOW].raygen.module = optix_module; + group_descs[PG_RGEN_SHADE_SHADOW].raygen.entryFunctionName = + "__raygen__kernel_optix_integrator_shade_shadow"; + group_descs[PG_RGEN_EVAL_DISPLACE].kind = OPTIX_PROGRAM_GROUP_KIND_RAYGEN; + group_descs[PG_RGEN_EVAL_DISPLACE].raygen.module = optix_module; + group_descs[PG_RGEN_EVAL_DISPLACE].raygen.entryFunctionName = + "__raygen__kernel_optix_shader_eval_displace"; + group_descs[PG_RGEN_EVAL_BACKGROUND].kind = OPTIX_PROGRAM_GROUP_KIND_RAYGEN; + group_descs[PG_RGEN_EVAL_BACKGROUND].raygen.module = optix_module; + group_descs[PG_RGEN_EVAL_BACKGROUND].raygen.entryFunctionName = + "__raygen__kernel_optix_shader_eval_background"; + group_descs[PG_RGEN_EVAL_CURVE_SHADOW_TRANSPARENCY].kind = OPTIX_PROGRAM_GROUP_KIND_RAYGEN; + group_descs[PG_RGEN_EVAL_CURVE_SHADOW_TRANSPARENCY].raygen.module = optix_module; + group_descs[PG_RGEN_EVAL_CURVE_SHADOW_TRANSPARENCY].raygen.entryFunctionName = + "__raygen__kernel_optix_shader_eval_curve_shadow_transparency"; + } + optix_assert(optixProgramGroupCreate( context, group_descs, NUM_PROGRAM_GROUPS, &group_options, nullptr, 0, groups)); @@ -680,7 +756,7 @@ bool OptiXDevice::load_kernels(const uint kernel_features) /* Set up SBT, which in this case is used only to select between different programs. */ sbt_data.alloc(NUM_PROGRAM_GROUPS); memset(sbt_data.host_pointer, 0, sizeof(SbtRecord) * NUM_PROGRAM_GROUPS); - for (unsigned int i = 0; i < NUM_PROGRAM_GROUPS; ++i) { + for (int i = 0; i < NUM_PROGRAM_GROUPS; ++i) { optix_assert(optixSbtRecordPackHeader(groups[i], &sbt_data[i])); optix_assert(optixProgramGroupGetStackSize(groups[i], &stack_size[i])); } @@ -704,25 +780,26 @@ bool OptiXDevice::load_kernels(const uint kernel_features) OptixPipelineLinkOptions link_options = {}; link_options.maxTraceDepth = 1; + link_options.debugLevel = module_options.debugLevel; - if (DebugFlags().optix.use_debug) { - link_options.debugLevel = OPTIX_COMPILE_DEBUG_LEVEL_FULL; - } - else { - link_options.debugLevel = OPTIX_COMPILE_DEBUG_LEVEL_NONE; - } - - if (kernel_features & KERNEL_FEATURE_NODE_RAYTRACE) { - /* Create shader raytracing pipeline. */ + if (kernel_features & (KERNEL_FEATURE_NODE_RAYTRACE | KERNEL_FEATURE_MNEE) && !use_osl) { + /* Create shader raytracing and MNEE pipeline. */ vector<OptixProgramGroup> pipeline_groups; pipeline_groups.reserve(NUM_PROGRAM_GROUPS); - pipeline_groups.push_back(groups[PG_RGEN_SHADE_SURFACE_RAYTRACE]); + if (kernel_features & KERNEL_FEATURE_NODE_RAYTRACE) { + pipeline_groups.push_back(groups[PG_RGEN_SHADE_SURFACE_RAYTRACE]); + pipeline_groups.push_back(groups[PG_CALL_SVM_AO]); + pipeline_groups.push_back(groups[PG_CALL_SVM_BEVEL]); + } + if (kernel_features & KERNEL_FEATURE_MNEE) { + pipeline_groups.push_back(groups[PG_RGEN_SHADE_SURFACE_MNEE]); + } pipeline_groups.push_back(groups[PG_MISS]); pipeline_groups.push_back(groups[PG_HITD]); pipeline_groups.push_back(groups[PG_HITS]); pipeline_groups.push_back(groups[PG_HITL]); pipeline_groups.push_back(groups[PG_HITV]); - if (motion_blur) { + if (pipeline_options.usesMotionBlur) { pipeline_groups.push_back(groups[PG_HITD_MOTION]); pipeline_groups.push_back(groups[PG_HITS_MOTION]); } @@ -730,8 +807,6 @@ bool OptiXDevice::load_kernels(const uint kernel_features) pipeline_groups.push_back(groups[PG_HITD_POINTCLOUD]); pipeline_groups.push_back(groups[PG_HITS_POINTCLOUD]); } - pipeline_groups.push_back(groups[PG_CALL_SVM_AO]); - pipeline_groups.push_back(groups[PG_CALL_SVM_BEVEL]); optix_assert(optixPipelineCreate(context, &pipeline_options, @@ -740,30 +815,33 @@ bool OptiXDevice::load_kernels(const uint kernel_features) pipeline_groups.size(), nullptr, 0, - &pipelines[PIP_SHADE_RAYTRACE])); + &pipelines[PIP_SHADE])); /* Combine ray generation and trace continuation stack size. */ - const unsigned int css = stack_size[PG_RGEN_SHADE_SURFACE_RAYTRACE].cssRG + + const unsigned int css = std::max(stack_size[PG_RGEN_SHADE_SURFACE_RAYTRACE].cssRG, + stack_size[PG_RGEN_SHADE_SURFACE_MNEE].cssRG) + link_options.maxTraceDepth * trace_css; const unsigned int dss = std::max(stack_size[PG_CALL_SVM_AO].dssDC, stack_size[PG_CALL_SVM_BEVEL].dssDC); /* Set stack size depending on pipeline options. */ optix_assert(optixPipelineSetStackSize( - pipelines[PIP_SHADE_RAYTRACE], 0, dss, css, motion_blur ? 3 : 2)); + pipelines[PIP_SHADE], 0, dss, css, pipeline_options.usesMotionBlur ? 3 : 2)); } - if (kernel_features & KERNEL_FEATURE_MNEE) { - /* Create MNEE pipeline. */ + { /* Create intersection-only pipeline. */ vector<OptixProgramGroup> pipeline_groups; pipeline_groups.reserve(NUM_PROGRAM_GROUPS); - pipeline_groups.push_back(groups[PG_RGEN_SHADE_SURFACE_MNEE]); + pipeline_groups.push_back(groups[PG_RGEN_INTERSECT_CLOSEST]); + pipeline_groups.push_back(groups[PG_RGEN_INTERSECT_SHADOW]); + pipeline_groups.push_back(groups[PG_RGEN_INTERSECT_SUBSURFACE]); + pipeline_groups.push_back(groups[PG_RGEN_INTERSECT_VOLUME_STACK]); pipeline_groups.push_back(groups[PG_MISS]); pipeline_groups.push_back(groups[PG_HITD]); pipeline_groups.push_back(groups[PG_HITS]); pipeline_groups.push_back(groups[PG_HITL]); pipeline_groups.push_back(groups[PG_HITV]); - if (motion_blur) { + if (pipeline_options.usesMotionBlur) { pipeline_groups.push_back(groups[PG_HITD_MOTION]); pipeline_groups.push_back(groups[PG_HITS_MOTION]); } @@ -771,8 +849,6 @@ bool OptiXDevice::load_kernels(const uint kernel_features) pipeline_groups.push_back(groups[PG_HITD_POINTCLOUD]); pipeline_groups.push_back(groups[PG_HITS_POINTCLOUD]); } - pipeline_groups.push_back(groups[PG_CALL_SVM_AO]); - pipeline_groups.push_back(groups[PG_CALL_SVM_BEVEL]); optix_assert(optixPipelineCreate(context, &pipeline_options, @@ -781,37 +857,234 @@ bool OptiXDevice::load_kernels(const uint kernel_features) pipeline_groups.size(), nullptr, 0, - &pipelines[PIP_SHADE_MNEE])); + &pipelines[PIP_INTERSECT])); - /* Combine ray generation and trace continuation stack size. */ - const unsigned int css = stack_size[PG_RGEN_SHADE_SURFACE_MNEE].cssRG + - link_options.maxTraceDepth * trace_css; - const unsigned int dss = 0; + /* Calculate continuation stack size based on the maximum of all ray generation stack sizes. */ + const unsigned int css = + std::max(stack_size[PG_RGEN_INTERSECT_CLOSEST].cssRG, + std::max(stack_size[PG_RGEN_INTERSECT_SHADOW].cssRG, + std::max(stack_size[PG_RGEN_INTERSECT_SUBSURFACE].cssRG, + stack_size[PG_RGEN_INTERSECT_VOLUME_STACK].cssRG))) + + link_options.maxTraceDepth * trace_css; - /* Set stack size depending on pipeline options. */ - optix_assert( - optixPipelineSetStackSize(pipelines[PIP_SHADE_MNEE], 0, dss, css, motion_blur ? 3 : 2)); + optix_assert(optixPipelineSetStackSize( + pipelines[PIP_INTERSECT], 0, 0, css, pipeline_options.usesMotionBlur ? 3 : 2)); } - { /* Create intersection-only pipeline. */ + return !have_error(); +} + +bool OptiXDevice::load_osl_kernels() +{ +# ifdef WITH_OSL + if (have_error()) { + return false; + } + + struct OSLKernel { + string ptx; + string init_entry; + string exec_entry; + }; + + /* This has to be in the same order as the ShaderType enum, so that the index calculation in + * osl_eval_nodes checks out */ + vector<OSLKernel> osl_kernels; + + for (ShaderType type = SHADER_TYPE_SURFACE; type <= SHADER_TYPE_BUMP; + type = static_cast<ShaderType>(type + 1)) { + const vector<OSL::ShaderGroupRef> &groups = (type == SHADER_TYPE_SURFACE ? + osl_globals.surface_state : + type == SHADER_TYPE_VOLUME ? + osl_globals.volume_state : + type == SHADER_TYPE_DISPLACEMENT ? + osl_globals.displacement_state : + osl_globals.bump_state); + for (const OSL::ShaderGroupRef &group : groups) { + if (group) { + string osl_ptx, init_name, entry_name; + osl_globals.ss->getattribute(group.get(), "group_init_name", init_name); + osl_globals.ss->getattribute(group.get(), "group_entry_name", entry_name); + osl_globals.ss->getattribute( + group.get(), "ptx_compiled_version", OSL::TypeDesc::PTR, &osl_ptx); + + int groupdata_size = 0; + osl_globals.ss->getattribute(group.get(), "groupdata_size", groupdata_size); + if (groupdata_size > 2048) { /* See 'group_data' array in kernel/osl/osl.h */ + set_error( + string_printf("Requested OSL group data size (%d) is greater than the maximum " + "supported with OptiX (2048)", + groupdata_size)); + return false; + } + + osl_kernels.push_back({std::move(osl_ptx), std::move(init_name), std::move(entry_name)}); + } + else { + /* Add empty entry for non-existent shader groups, so that the index stays stable. */ + osl_kernels.emplace_back(); + } + } + } + + const CUDAContextScope scope(this); + + if (pipelines[PIP_SHADE]) { + optixPipelineDestroy(pipelines[PIP_SHADE]); + } + + for (OptixModule &module : osl_modules) { + if (module != NULL) { + optixModuleDestroy(module); + module = NULL; + } + } + for (OptixProgramGroup &group : osl_groups) { + if (group != NULL) { + optixProgramGroupDestroy(group); + group = NULL; + } + } + + OptixProgramGroupOptions group_options = {}; /* There are no options currently. */ + OptixModuleCompileOptions module_options = {}; + module_options.optLevel = OPTIX_COMPILE_OPTIMIZATION_LEVEL_3; + module_options.debugLevel = OPTIX_COMPILE_DEBUG_LEVEL_NONE; + + osl_groups.resize(osl_kernels.size() * 2 + 1); + osl_modules.resize(osl_kernels.size() + 1); + + { /* Load and compile PTX module with OSL services. */ + string ptx_data, ptx_filename = path_get("lib/kernel_optix_osl_services.ptx"); + if (!path_read_text(ptx_filename, ptx_data)) { + set_error(string_printf("Failed to load OptiX OSL services kernel from '%s'", + ptx_filename.c_str())); + return false; + } + + const OptixResult result = optixModuleCreateFromPTX(context, + &module_options, + &pipeline_options, + ptx_data.data(), + ptx_data.size(), + nullptr, + 0, + &osl_modules.back()); + if (result != OPTIX_SUCCESS) { + set_error(string_printf("Failed to load OptiX OSL services kernel from '%s' (%s)", + ptx_filename.c_str(), + optixGetErrorName(result))); + return false; + } + + OptixProgramGroupDesc group_desc = {}; + group_desc.kind = OPTIX_PROGRAM_GROUP_KIND_CALLABLES; + group_desc.callables.entryFunctionNameDC = "__direct_callable__dummy_services"; + group_desc.callables.moduleDC = osl_modules.back(); + + optix_assert(optixProgramGroupCreate( + context, &group_desc, 1, &group_options, nullptr, 0, &osl_groups.back())); + } + + TaskPool pool; + vector<OptixResult> results(osl_kernels.size(), OPTIX_SUCCESS); + + for (size_t i = 0; i < osl_kernels.size(); ++i) { + if (osl_kernels[i].ptx.empty()) { + continue; + } + +# if OPTIX_ABI_VERSION >= 55 + OptixTask task = nullptr; + results[i] = optixModuleCreateFromPTXWithTasks(context, + &module_options, + &pipeline_options, + osl_kernels[i].ptx.data(), + osl_kernels[i].ptx.size(), + nullptr, + nullptr, + &osl_modules[i], + &task); + if (results[i] == OPTIX_SUCCESS) { + execute_optix_task(pool, task, results[i]); + } +# else + pool.push([this, &results, i, &module_options, &osl_kernels]() { + results[i] = optixModuleCreateFromPTX(context, + &module_options, + &pipeline_options, + osl_kernels[i].ptx.data(), + osl_kernels[i].ptx.size(), + nullptr, + 0, + &osl_modules[i]); + }); +# endif + } + + pool.wait_work(); + + for (size_t i = 0; i < osl_kernels.size(); ++i) { + if (osl_kernels[i].ptx.empty()) { + continue; + } + + if (results[i] != OPTIX_SUCCESS) { + set_error(string_printf("Failed to load OptiX OSL kernel for %s (%s)", + osl_kernels[i].init_entry.c_str(), + optixGetErrorName(results[i]))); + return false; + } + + OptixProgramGroupDesc group_descs[2] = {}; + group_descs[0].kind = OPTIX_PROGRAM_GROUP_KIND_CALLABLES; + group_descs[0].callables.entryFunctionNameDC = osl_kernels[i].init_entry.c_str(); + group_descs[0].callables.moduleDC = osl_modules[i]; + group_descs[1].kind = OPTIX_PROGRAM_GROUP_KIND_CALLABLES; + group_descs[1].callables.entryFunctionNameDC = osl_kernels[i].exec_entry.c_str(); + group_descs[1].callables.moduleDC = osl_modules[i]; + + optix_assert(optixProgramGroupCreate( + context, group_descs, 2, &group_options, nullptr, 0, &osl_groups[i * 2])); + } + + vector<OptixStackSizes> osl_stack_size(osl_groups.size()); + + /* Update SBT with new entries. */ + sbt_data.alloc(NUM_PROGRAM_GROUPS + osl_groups.size()); + for (int i = 0; i < NUM_PROGRAM_GROUPS; ++i) { + optix_assert(optixSbtRecordPackHeader(groups[i], &sbt_data[i])); + } + for (size_t i = 0; i < osl_groups.size(); ++i) { + if (osl_groups[i] != NULL) { + optix_assert(optixSbtRecordPackHeader(osl_groups[i], &sbt_data[NUM_PROGRAM_GROUPS + i])); + optix_assert(optixProgramGroupGetStackSize(osl_groups[i], &osl_stack_size[i])); + } + } + sbt_data.copy_to_device(); /* Upload updated SBT to device. */ + + OptixPipelineLinkOptions link_options = {}; + link_options.maxTraceDepth = 0; + link_options.debugLevel = OPTIX_COMPILE_DEBUG_LEVEL_NONE; + + { vector<OptixProgramGroup> pipeline_groups; pipeline_groups.reserve(NUM_PROGRAM_GROUPS); - pipeline_groups.push_back(groups[PG_RGEN_INTERSECT_CLOSEST]); - pipeline_groups.push_back(groups[PG_RGEN_INTERSECT_SHADOW]); - pipeline_groups.push_back(groups[PG_RGEN_INTERSECT_SUBSURFACE]); - pipeline_groups.push_back(groups[PG_RGEN_INTERSECT_VOLUME_STACK]); - pipeline_groups.push_back(groups[PG_MISS]); - pipeline_groups.push_back(groups[PG_HITD]); - pipeline_groups.push_back(groups[PG_HITS]); - pipeline_groups.push_back(groups[PG_HITL]); - pipeline_groups.push_back(groups[PG_HITV]); - if (motion_blur) { - pipeline_groups.push_back(groups[PG_HITD_MOTION]); - pipeline_groups.push_back(groups[PG_HITS_MOTION]); - } - if (kernel_features & KERNEL_FEATURE_POINTCLOUD) { - pipeline_groups.push_back(groups[PG_HITD_POINTCLOUD]); - pipeline_groups.push_back(groups[PG_HITS_POINTCLOUD]); + pipeline_groups.push_back(groups[PG_RGEN_SHADE_BACKGROUND]); + pipeline_groups.push_back(groups[PG_RGEN_SHADE_LIGHT]); + pipeline_groups.push_back(groups[PG_RGEN_SHADE_SURFACE]); + pipeline_groups.push_back(groups[PG_RGEN_SHADE_SURFACE_RAYTRACE]); + pipeline_groups.push_back(groups[PG_RGEN_SHADE_SURFACE_MNEE]); + pipeline_groups.push_back(groups[PG_RGEN_SHADE_VOLUME]); + pipeline_groups.push_back(groups[PG_RGEN_SHADE_SHADOW]); + pipeline_groups.push_back(groups[PG_RGEN_EVAL_DISPLACE]); + pipeline_groups.push_back(groups[PG_RGEN_EVAL_BACKGROUND]); + pipeline_groups.push_back(groups[PG_RGEN_EVAL_CURVE_SHADOW_TRANSPARENCY]); + + for (const OptixProgramGroup &group : osl_groups) { + if (group != NULL) { + pipeline_groups.push_back(group); + } } optix_assert(optixPipelineCreate(context, @@ -821,26 +1094,30 @@ bool OptiXDevice::load_kernels(const uint kernel_features) pipeline_groups.size(), nullptr, 0, - &pipelines[PIP_INTERSECT])); + &pipelines[PIP_SHADE])); - /* Calculate continuation stack size based on the maximum of all ray generation stack sizes. */ - const unsigned int css = - std::max(stack_size[PG_RGEN_INTERSECT_CLOSEST].cssRG, - std::max(stack_size[PG_RGEN_INTERSECT_SHADOW].cssRG, - std::max(stack_size[PG_RGEN_INTERSECT_SUBSURFACE].cssRG, - stack_size[PG_RGEN_INTERSECT_VOLUME_STACK].cssRG))) + - link_options.maxTraceDepth * trace_css; + unsigned int dss = 0; + for (unsigned int i = 0; i < osl_stack_size.size(); ++i) { + dss = std::max(dss, osl_stack_size[i].dssDC); + } - optix_assert( - optixPipelineSetStackSize(pipelines[PIP_INTERSECT], 0, 0, css, motion_blur ? 3 : 2)); + optix_assert(optixPipelineSetStackSize( + pipelines[PIP_SHADE], 0, dss, 0, pipeline_options.usesMotionBlur ? 3 : 2)); } - /* Clean up program group objects. */ - for (unsigned int i = 0; i < NUM_PROGRAM_GROUPS; ++i) { - optixProgramGroupDestroy(groups[i]); - } + return !have_error(); +# else + return false; +# endif +} - return true; +void *OptiXDevice::get_cpu_osl_memory() +{ +# ifdef WITH_OSL + return &osl_globals; +# else + return NULL; +# endif } /* -------------------------------------------------------------------- @@ -1567,7 +1844,7 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit) size_t num_motion_steps = 1; Attribute *motion_keys = hair->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION); - if (motion_blur && hair->get_use_motion_blur() && motion_keys) { + if (pipeline_options.usesMotionBlur && hair->get_use_motion_blur() && motion_keys) { num_motion_steps = hair->get_motion_steps(); } @@ -1721,7 +1998,7 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit) size_t num_motion_steps = 1; Attribute *motion_keys = mesh->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION); - if (motion_blur && mesh->get_use_motion_blur() && motion_keys) { + if (pipeline_options.usesMotionBlur && mesh->get_use_motion_blur() && motion_keys) { num_motion_steps = mesh->get_motion_steps(); } @@ -1788,7 +2065,7 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit) size_t num_motion_steps = 1; Attribute *motion_points = pointcloud->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION); - if (motion_blur && pointcloud->get_use_motion_blur() && motion_points) { + if (pipeline_options.usesMotionBlur && pointcloud->get_use_motion_blur() && motion_points) { num_motion_steps = pointcloud->get_motion_steps(); } @@ -1885,7 +2162,7 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit) /* Calculate total motion transform size and allocate memory for them. */ size_t motion_transform_offset = 0; - if (motion_blur) { + if (pipeline_options.usesMotionBlur) { size_t total_motion_transform_size = 0; for (Object *const ob : bvh->objects) { if (ob->is_traceable() && ob->use_motion()) { @@ -1936,7 +2213,7 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit) if (ob->get_geometry()->geometry_type == Geometry::HAIR && static_cast<const Hair *>(ob->get_geometry())->curve_shape == CURVE_THICK) { - if (motion_blur && ob->get_geometry()->has_motion_blur()) { + if (pipeline_options.usesMotionBlur && ob->get_geometry()->has_motion_blur()) { /* Select between motion blur and non-motion blur built-in intersection module. */ instance.sbtOffset = PG_HITD_MOTION - PG_HITD; } @@ -1964,7 +2241,7 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit) } /* Insert motion traversable if object has motion. */ - if (motion_blur && ob->use_motion()) { + if (pipeline_options.usesMotionBlur && ob->use_motion()) { size_t motion_keys = max(ob->get_motion().size(), (size_t)2) - 2; size_t motion_transform_size = sizeof(OptixSRTMotionTransform) + motion_keys * sizeof(OptixSRTData); diff --git a/intern/cycles/device/optix/device_impl.h b/intern/cycles/device/optix/device_impl.h index 76c8af9bc3f..ad0e7b93454 100644 --- a/intern/cycles/device/optix/device_impl.h +++ b/intern/cycles/device/optix/device_impl.h @@ -9,6 +9,7 @@ # include "device/cuda/device_impl.h" # include "device/optix/queue.h" # include "device/optix/util.h" +# include "kernel/osl/globals.h" # include "kernel/types.h" # include "util/unique_ptr.h" @@ -23,8 +24,16 @@ enum { PG_RGEN_INTERSECT_SHADOW, PG_RGEN_INTERSECT_SUBSURFACE, PG_RGEN_INTERSECT_VOLUME_STACK, + PG_RGEN_SHADE_BACKGROUND, + PG_RGEN_SHADE_LIGHT, + PG_RGEN_SHADE_SURFACE, PG_RGEN_SHADE_SURFACE_RAYTRACE, PG_RGEN_SHADE_SURFACE_MNEE, + PG_RGEN_SHADE_VOLUME, + PG_RGEN_SHADE_SHADOW, + PG_RGEN_EVAL_DISPLACE, + PG_RGEN_EVAL_BACKGROUND, + PG_RGEN_EVAL_CURVE_SHADOW_TRANSPARENCY, PG_MISS, PG_HITD, /* Default hit group. */ PG_HITS, /* __SHADOW_RECORD_ALL__ hit group. */ @@ -40,14 +49,14 @@ enum { }; static const int MISS_PROGRAM_GROUP_OFFSET = PG_MISS; -static const int NUM_MIS_PROGRAM_GROUPS = 1; +static const int NUM_MISS_PROGRAM_GROUPS = 1; static const int HIT_PROGAM_GROUP_OFFSET = PG_HITD; static const int NUM_HIT_PROGRAM_GROUPS = 8; static const int CALLABLE_PROGRAM_GROUPS_BASE = PG_CALL_SVM_AO; static const int NUM_CALLABLE_PROGRAM_GROUPS = 2; /* List of OptiX pipelines. */ -enum { PIP_SHADE_RAYTRACE, PIP_SHADE_MNEE, PIP_INTERSECT, NUM_PIPELINES }; +enum { PIP_SHADE, PIP_INTERSECT, NUM_PIPELINES }; /* A single shader binding table entry. */ struct SbtRecord { @@ -61,12 +70,20 @@ class OptiXDevice : public CUDADevice { OptixModule optix_module = NULL; /* All necessary OptiX kernels are in one module. */ OptixModule builtin_modules[2] = {}; OptixPipeline pipelines[NUM_PIPELINES] = {}; + OptixProgramGroup groups[NUM_PROGRAM_GROUPS] = {}; + OptixPipelineCompileOptions pipeline_options = {}; - bool motion_blur = false; device_vector<SbtRecord> sbt_data; device_only_memory<KernelParamsOptiX> launch_params; - OptixTraversableHandle tlas_handle = 0; +# ifdef WITH_OSL + OSLGlobals osl_globals; + vector<OptixModule> osl_modules; + vector<OptixProgramGroup> osl_groups; +# endif + + private: + OptixTraversableHandle tlas_handle = 0; vector<unique_ptr<device_only_memory<char>>> delayed_free_bvh_memory; thread_mutex delayed_free_bvh_mutex; @@ -100,13 +117,14 @@ class OptiXDevice : public CUDADevice { OptiXDevice(const DeviceInfo &info, Stats &stats, Profiler &profiler); ~OptiXDevice(); - private: BVHLayoutMask get_bvh_layout_mask() const override; string compile_kernel_get_common_cflags(const uint kernel_features); bool load_kernels(const uint kernel_features) override; + bool load_osl_kernels() override; + bool build_optix_bvh(BVHOptiX *bvh, OptixBuildOperation operation, const OptixBuildInput &build_input, @@ -123,6 +141,8 @@ class OptiXDevice : public CUDADevice { virtual unique_ptr<DeviceQueue> gpu_queue_create() override; + void *get_cpu_osl_memory() override; + /* -------------------------------------------------------------------- * Denoising. */ diff --git a/intern/cycles/device/optix/queue.cpp b/intern/cycles/device/optix/queue.cpp index 3bc547ed11d..1bfd154d449 100644 --- a/intern/cycles/device/optix/queue.cpp +++ b/intern/cycles/device/optix/queue.cpp @@ -24,21 +24,33 @@ void OptiXDeviceQueue::init_execution() CUDADeviceQueue::init_execution(); } -static bool is_optix_specific_kernel(DeviceKernel kernel) +static bool is_optix_specific_kernel(DeviceKernel kernel, bool use_osl) { - return (kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE || - kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_MNEE || - kernel == DEVICE_KERNEL_INTEGRATOR_INTERSECT_CLOSEST || - kernel == DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW || - kernel == DEVICE_KERNEL_INTEGRATOR_INTERSECT_SUBSURFACE || - kernel == DEVICE_KERNEL_INTEGRATOR_INTERSECT_VOLUME_STACK); +# ifdef WITH_OSL + /* OSL uses direct callables to execute, so shading needs to be done in OptiX if OSL is used. */ + if (use_osl && device_kernel_has_shading(kernel)) { + return true; + } +# else + (void)use_osl; +# endif + + return device_kernel_has_intersection(kernel); } bool OptiXDeviceQueue::enqueue(DeviceKernel kernel, const int work_size, DeviceKernelArguments const &args) { - if (!is_optix_specific_kernel(kernel)) { + OptiXDevice *const optix_device = static_cast<OptiXDevice *>(cuda_device_); + +# ifdef WITH_OSL + const bool use_osl = static_cast<OSLGlobals *>(optix_device->get_cpu_osl_memory())->use; +# else + const bool use_osl = false; +# endif + + if (!is_optix_specific_kernel(kernel, use_osl)) { return CUDADeviceQueue::enqueue(kernel, work_size, args); } @@ -50,8 +62,6 @@ bool OptiXDeviceQueue::enqueue(DeviceKernel kernel, const CUDAContextScope scope(cuda_device_); - OptiXDevice *const optix_device = static_cast<OptiXDevice *>(cuda_device_); - const device_ptr sbt_data_ptr = optix_device->sbt_data.device_pointer; const device_ptr launch_params_ptr = optix_device->launch_params.device_pointer; @@ -62,9 +72,7 @@ bool OptiXDeviceQueue::enqueue(DeviceKernel kernel, sizeof(device_ptr), cuda_stream_)); - if (kernel == DEVICE_KERNEL_INTEGRATOR_INTERSECT_CLOSEST || - kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE || - kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_MNEE) { + if (kernel == DEVICE_KERNEL_INTEGRATOR_INTERSECT_CLOSEST || device_kernel_has_shading(kernel)) { cuda_device_assert( cuda_device_, cuMemcpyHtoDAsync(launch_params_ptr + offsetof(KernelParamsOptiX, render_buffer), @@ -72,6 +80,15 @@ bool OptiXDeviceQueue::enqueue(DeviceKernel kernel, sizeof(device_ptr), cuda_stream_)); } + if (kernel == DEVICE_KERNEL_SHADER_EVAL_DISPLACE || + kernel == DEVICE_KERNEL_SHADER_EVAL_BACKGROUND || + kernel == DEVICE_KERNEL_SHADER_EVAL_CURVE_SHADOW_TRANSPARENCY) { + cuda_device_assert(cuda_device_, + cuMemcpyHtoDAsync(launch_params_ptr + offsetof(KernelParamsOptiX, offset), + args.values[2], // &d_offset + sizeof(int32_t), + cuda_stream_)); + } cuda_device_assert(cuda_device_, cuStreamSynchronize(cuda_stream_)); @@ -79,14 +96,35 @@ bool OptiXDeviceQueue::enqueue(DeviceKernel kernel, OptixShaderBindingTable sbt_params = {}; switch (kernel) { + case DEVICE_KERNEL_INTEGRATOR_SHADE_BACKGROUND: + pipeline = optix_device->pipelines[PIP_SHADE]; + sbt_params.raygenRecord = sbt_data_ptr + PG_RGEN_SHADE_BACKGROUND * sizeof(SbtRecord); + break; + case DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT: + pipeline = optix_device->pipelines[PIP_SHADE]; + sbt_params.raygenRecord = sbt_data_ptr + PG_RGEN_SHADE_LIGHT * sizeof(SbtRecord); + break; + case DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE: + pipeline = optix_device->pipelines[PIP_SHADE]; + sbt_params.raygenRecord = sbt_data_ptr + PG_RGEN_SHADE_SURFACE * sizeof(SbtRecord); + break; case DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE: - pipeline = optix_device->pipelines[PIP_SHADE_RAYTRACE]; + pipeline = optix_device->pipelines[PIP_SHADE]; sbt_params.raygenRecord = sbt_data_ptr + PG_RGEN_SHADE_SURFACE_RAYTRACE * sizeof(SbtRecord); break; case DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_MNEE: - pipeline = optix_device->pipelines[PIP_SHADE_MNEE]; + pipeline = optix_device->pipelines[PIP_SHADE]; sbt_params.raygenRecord = sbt_data_ptr + PG_RGEN_SHADE_SURFACE_MNEE * sizeof(SbtRecord); break; + case DEVICE_KERNEL_INTEGRATOR_SHADE_VOLUME: + pipeline = optix_device->pipelines[PIP_SHADE]; + sbt_params.raygenRecord = sbt_data_ptr + PG_RGEN_SHADE_VOLUME * sizeof(SbtRecord); + break; + case DEVICE_KERNEL_INTEGRATOR_SHADE_SHADOW: + pipeline = optix_device->pipelines[PIP_SHADE]; + sbt_params.raygenRecord = sbt_data_ptr + PG_RGEN_SHADE_SHADOW * sizeof(SbtRecord); + break; + case DEVICE_KERNEL_INTEGRATOR_INTERSECT_CLOSEST: pipeline = optix_device->pipelines[PIP_INTERSECT]; sbt_params.raygenRecord = sbt_data_ptr + PG_RGEN_INTERSECT_CLOSEST * sizeof(SbtRecord); @@ -104,6 +142,20 @@ bool OptiXDeviceQueue::enqueue(DeviceKernel kernel, sbt_params.raygenRecord = sbt_data_ptr + PG_RGEN_INTERSECT_VOLUME_STACK * sizeof(SbtRecord); break; + case DEVICE_KERNEL_SHADER_EVAL_DISPLACE: + pipeline = optix_device->pipelines[PIP_SHADE]; + sbt_params.raygenRecord = sbt_data_ptr + PG_RGEN_EVAL_DISPLACE * sizeof(SbtRecord); + break; + case DEVICE_KERNEL_SHADER_EVAL_BACKGROUND: + pipeline = optix_device->pipelines[PIP_SHADE]; + sbt_params.raygenRecord = sbt_data_ptr + PG_RGEN_EVAL_BACKGROUND * sizeof(SbtRecord); + break; + case DEVICE_KERNEL_SHADER_EVAL_CURVE_SHADOW_TRANSPARENCY: + pipeline = optix_device->pipelines[PIP_SHADE]; + sbt_params.raygenRecord = sbt_data_ptr + + PG_RGEN_EVAL_CURVE_SHADOW_TRANSPARENCY * sizeof(SbtRecord); + break; + default: LOG(ERROR) << "Invalid kernel " << device_kernel_as_string(kernel) << " is attempted to be enqueued."; @@ -112,7 +164,7 @@ bool OptiXDeviceQueue::enqueue(DeviceKernel kernel, sbt_params.missRecordBase = sbt_data_ptr + MISS_PROGRAM_GROUP_OFFSET * sizeof(SbtRecord); sbt_params.missRecordStrideInBytes = sizeof(SbtRecord); - sbt_params.missRecordCount = NUM_MIS_PROGRAM_GROUPS; + sbt_params.missRecordCount = NUM_MISS_PROGRAM_GROUPS; sbt_params.hitgroupRecordBase = sbt_data_ptr + HIT_PROGAM_GROUP_OFFSET * sizeof(SbtRecord); sbt_params.hitgroupRecordStrideInBytes = sizeof(SbtRecord); sbt_params.hitgroupRecordCount = NUM_HIT_PROGRAM_GROUPS; @@ -120,6 +172,12 @@ bool OptiXDeviceQueue::enqueue(DeviceKernel kernel, sbt_params.callablesRecordCount = NUM_CALLABLE_PROGRAM_GROUPS; sbt_params.callablesRecordStrideInBytes = sizeof(SbtRecord); +# ifdef WITH_OSL + if (use_osl) { + sbt_params.callablesRecordCount += static_cast<unsigned int>(optix_device->osl_groups.size()); + } +# endif + /* Launch the ray generation program. */ optix_device_assert(optix_device, optixLaunch(pipeline, diff --git a/intern/cycles/kernel/CMakeLists.txt b/intern/cycles/kernel/CMakeLists.txt index 3fbb346e94f..99f9e536977 100644 --- a/intern/cycles/kernel/CMakeLists.txt +++ b/intern/cycles/kernel/CMakeLists.txt @@ -37,6 +37,14 @@ set(SRC_KERNEL_DEVICE_OPTIX device/optix/kernel_shader_raytrace.cu ) +if(WITH_CYCLES_OSL AND (OSL_LIBRARY_VERSION_MINOR GREATER_EQUAL 13 OR OSL_LIBRARY_VERSION_MAJOR GREATER 1)) + set(SRC_KERNEL_DEVICE_OPTIX + ${SRC_KERNEL_DEVICE_OPTIX} + osl/services_optix.cu + device/optix/kernel_osl.cu + ) +endif() + set(SRC_KERNEL_DEVICE_ONEAPI device/oneapi/kernel.cpp ) @@ -181,6 +189,16 @@ set(SRC_KERNEL_SVM_HEADERS svm/vertex_color.h ) +if(WITH_CYCLES_OSL) + set(SRC_KERNEL_OSL_HEADERS + osl/osl.h + osl/closures_setup.h + osl/closures_template.h + osl/services_gpu.h + osl/types.h + ) +endif() + set(SRC_KERNEL_GEOM_HEADERS geom/geom.h geom/attribute.h @@ -306,6 +324,7 @@ set(SRC_KERNEL_HEADERS ${SRC_KERNEL_GEOM_HEADERS} ${SRC_KERNEL_INTEGRATOR_HEADERS} ${SRC_KERNEL_LIGHT_HEADERS} + ${SRC_KERNEL_OSL_HEADERS} ${SRC_KERNEL_SAMPLE_HEADERS} ${SRC_KERNEL_SVM_HEADERS} ${SRC_KERNEL_TYPES_HEADERS} @@ -708,6 +727,16 @@ if(WITH_CYCLES_DEVICE_OPTIX AND WITH_CYCLES_CUDA_BINARIES) kernel_optix_shader_raytrace "device/optix/kernel_shader_raytrace.cu" "--keep-device-functions") + if(WITH_CYCLES_OSL AND (OSL_LIBRARY_VERSION_MINOR GREATER_EQUAL 13 OR OSL_LIBRARY_VERSION_MAJOR GREATER 1)) + CYCLES_OPTIX_KERNEL_ADD( + kernel_optix_osl + "device/optix/kernel_osl.cu" + "--relocatable-device-code=true") + CYCLES_OPTIX_KERNEL_ADD( + kernel_optix_osl_services + "osl/services_optix.cu" + "--relocatable-device-code=true") + endif() add_custom_target(cycles_kernel_optix ALL DEPENDS ${optix_ptx}) cycles_set_solution_folder(cycles_kernel_optix) @@ -995,6 +1024,7 @@ source_group("geom" FILES ${SRC_KERNEL_GEOM_HEADERS}) source_group("integrator" FILES ${SRC_KERNEL_INTEGRATOR_HEADERS}) source_group("kernel" FILES ${SRC_KERNEL_TYPES_HEADERS}) source_group("light" FILES ${SRC_KERNEL_LIGHT_HEADERS}) +source_group("osl" FILES ${SRC_KERNEL_OSL_HEADERS}) source_group("sample" FILES ${SRC_KERNEL_SAMPLE_HEADERS}) source_group("svm" FILES ${SRC_KERNEL_SVM_HEADERS}) source_group("util" FILES ${SRC_KERNEL_UTIL_HEADERS}) @@ -1031,6 +1061,7 @@ delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_KERNEL_FILM_HEADERS}" ${CYCLE delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_KERNEL_GEOM_HEADERS}" ${CYCLES_INSTALL_PATH}/source/kernel/geom) delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_KERNEL_INTEGRATOR_HEADERS}" ${CYCLES_INSTALL_PATH}/source/kernel/integrator) delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_KERNEL_LIGHT_HEADERS}" ${CYCLES_INSTALL_PATH}/source/kernel/light) +delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_KERNEL_OSL_HEADERS}" ${CYCLES_INSTALL_PATH}/source/kernel/osl) delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_KERNEL_SAMPLE_HEADERS}" ${CYCLES_INSTALL_PATH}/source/kernel/sample) delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_KERNEL_SVM_HEADERS}" ${CYCLES_INSTALL_PATH}/source/kernel/svm) delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_KERNEL_TYPES_HEADERS}" ${CYCLES_INSTALL_PATH}/source/kernel) diff --git a/intern/cycles/kernel/closure/bsdf.h b/intern/cycles/kernel/closure/bsdf.h index 71af68aa80e..2f5c5d7bd0c 100644 --- a/intern/cycles/kernel/closure/bsdf.h +++ b/intern/cycles/kernel/closure/bsdf.h @@ -297,8 +297,10 @@ ccl_device_inline void bsdf_roughness_eta(const KernelGlobals kg, ccl_private float2 *roughness, ccl_private float *eta) { +#ifdef __SVM__ bool refractive = false; float alpha = 1.0f; +#endif switch (sc->type) { case CLOSURE_BSDF_DIFFUSE_ID: *roughness = one_float2(); diff --git a/intern/cycles/kernel/device/cuda/compat.h b/intern/cycles/kernel/device/cuda/compat.h index 51e1381d552..3a950779c11 100644 --- a/intern/cycles/kernel/device/cuda/compat.h +++ b/intern/cycles/kernel/device/cuda/compat.h @@ -30,6 +30,7 @@ typedef unsigned long long uint64_t; /* Qualifiers */ #define ccl_device __device__ __inline__ +#define ccl_device_extern extern "C" __device__ #if __CUDA_ARCH__ < 500 # define ccl_device_inline __device__ __forceinline__ # define ccl_device_forceinline __device__ __forceinline__ @@ -109,14 +110,14 @@ ccl_device_forceinline T ccl_gpu_tex_object_read_3D(const ccl_gpu_tex_object_3D typedef unsigned short half; -__device__ half __float2half(const float f) +ccl_device_forceinline half __float2half(const float f) { half val; asm("{ cvt.rn.f16.f32 %0, %1;}\n" : "=h"(val) : "f"(f)); return val; } -__device__ float __half2float(const half h) +ccl_device_forceinline float __half2float(const half h) { float val; asm("{ cvt.f32.f16 %0, %1;}\n" : "=f"(val) : "h"(h)); diff --git a/intern/cycles/kernel/device/hip/compat.h b/intern/cycles/kernel/device/hip/compat.h index 648988c31b6..8755395c82c 100644 --- a/intern/cycles/kernel/device/hip/compat.h +++ b/intern/cycles/kernel/device/hip/compat.h @@ -28,6 +28,7 @@ typedef unsigned long long uint64_t; /* Qualifiers */ #define ccl_device __device__ __inline__ +#define ccl_device_extern extern "C" __device__ #define ccl_device_inline __device__ __inline__ #define ccl_device_forceinline __device__ __forceinline__ #define ccl_device_noinline __device__ __noinline__ diff --git a/intern/cycles/kernel/device/metal/compat.h b/intern/cycles/kernel/device/metal/compat.h index f689e93e5a2..2dd6cc98b59 100644 --- a/intern/cycles/kernel/device/metal/compat.h +++ b/intern/cycles/kernel/device/metal/compat.h @@ -38,6 +38,7 @@ using namespace metal::raytracing; # define ccl_device_noinline ccl_device __attribute__((noinline)) #endif +#define ccl_device_extern extern "C" #define ccl_device_noinline_cpu ccl_device #define ccl_device_inline_method ccl_device #define ccl_global device diff --git a/intern/cycles/kernel/device/oneapi/compat.h b/intern/cycles/kernel/device/oneapi/compat.h index dfaec65130c..b83512180d7 100644 --- a/intern/cycles/kernel/device/oneapi/compat.h +++ b/intern/cycles/kernel/device/oneapi/compat.h @@ -28,6 +28,7 @@ /* Qualifier wrappers for different names on different devices */ #define ccl_device +#define ccl_device_extern extern "C" #define ccl_global #define ccl_always_inline __attribute__((always_inline)) #define ccl_device_inline inline diff --git a/intern/cycles/kernel/device/optix/compat.h b/intern/cycles/kernel/device/optix/compat.h index 1a11a533b7e..e13101f57b8 100644 --- a/intern/cycles/kernel/device/optix/compat.h +++ b/intern/cycles/kernel/device/optix/compat.h @@ -33,14 +33,16 @@ typedef unsigned long long uint64_t; #endif #define ccl_device \ - __device__ __forceinline__ // Function calls are bad for OptiX performance, so inline everything + static __device__ \ + __forceinline__ // Function calls are bad for OptiX performance, so inline everything +#define ccl_device_extern extern "C" __device__ #define ccl_device_inline ccl_device #define ccl_device_forceinline ccl_device -#define ccl_device_inline_method ccl_device -#define ccl_device_noinline __device__ __noinline__ +#define ccl_device_inline_method __device__ __forceinline__ +#define ccl_device_noinline static __device__ __noinline__ #define ccl_device_noinline_cpu ccl_device #define ccl_global -#define ccl_inline_constant __constant__ +#define ccl_inline_constant static __constant__ #define ccl_device_constant __constant__ __device__ #define ccl_constant const #define ccl_gpu_shared __shared__ @@ -57,23 +59,6 @@ typedef unsigned long long uint64_t; #define kernel_assert(cond) -/* GPU thread, block, grid size and index */ - -#define ccl_gpu_thread_idx_x (threadIdx.x) -#define ccl_gpu_block_dim_x (blockDim.x) -#define ccl_gpu_block_idx_x (blockIdx.x) -#define ccl_gpu_grid_dim_x (gridDim.x) -#define ccl_gpu_warp_size (warpSize) -#define ccl_gpu_thread_mask(thread_warp) uint(0xFFFFFFFF >> (ccl_gpu_warp_size - thread_warp)) - -#define ccl_gpu_global_id_x() (ccl_gpu_block_idx_x * ccl_gpu_block_dim_x + ccl_gpu_thread_idx_x) -#define ccl_gpu_global_size_x() (ccl_gpu_grid_dim_x * ccl_gpu_block_dim_x) - -/* GPU warp synchronization. */ - -#define ccl_gpu_syncthreads() __syncthreads() -#define ccl_gpu_ballot(predicate) __ballot_sync(0xFFFFFFFF, predicate) - /* GPU texture objects */ typedef unsigned long long CUtexObject; @@ -101,14 +86,14 @@ ccl_device_forceinline T ccl_gpu_tex_object_read_3D(const ccl_gpu_tex_object_3D typedef unsigned short half; -__device__ half __float2half(const float f) +ccl_device_forceinline half __float2half(const float f) { half val; asm("{ cvt.rn.f16.f32 %0, %1;}\n" : "=h"(val) : "f"(f)); return val; } -__device__ float __half2float(const half h) +ccl_device_forceinline float __half2float(const half h) { float val; asm("{ cvt.f32.f16 %0, %1;}\n" : "=f"(val) : "h"(h)); diff --git a/intern/cycles/kernel/device/optix/globals.h b/intern/cycles/kernel/device/optix/globals.h index 7af2e421378..126df74bc8c 100644 --- a/intern/cycles/kernel/device/optix/globals.h +++ b/intern/cycles/kernel/device/optix/globals.h @@ -25,6 +25,7 @@ struct KernelParamsOptiX { /* Kernel arguments */ const int *path_index_array; float *render_buffer; + int offset; /* Global scene data and textures */ KernelData data; @@ -36,7 +37,11 @@ struct KernelParamsOptiX { }; #ifdef __NVCC__ -extern "C" static __constant__ KernelParamsOptiX kernel_params; +extern "C" +# ifndef __CUDACC_RDC__ + static +# endif + __constant__ KernelParamsOptiX kernel_params; #endif /* Abstraction macros */ diff --git a/intern/cycles/kernel/device/optix/kernel_osl.cu b/intern/cycles/kernel/device/optix/kernel_osl.cu new file mode 100644 index 00000000000..0f3f477935b --- /dev/null +++ b/intern/cycles/kernel/device/optix/kernel_osl.cu @@ -0,0 +1,83 @@ +/* SPDX-License-Identifier: Apache-2.0 + * Copyright 2011-2022 Blender Foundation */ + +#define WITH_OSL + +/* Copy of the regular OptiX kernels with additional OSL support. */ + +#include "kernel/device/optix/kernel_shader_raytrace.cu" + +#include "kernel/bake/bake.h" +#include "kernel/integrator/shade_background.h" +#include "kernel/integrator/shade_light.h" +#include "kernel/integrator/shade_shadow.h" +#include "kernel/integrator/shade_volume.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() +{ + 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(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_volume() +{ + 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_volume(nullptr, path_index, kernel_params.render_buffer); +} + +extern "C" __global__ void __raygen__kernel_optix_integrator_shade_shadow() +{ + 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_shadow(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; + const int global_index = kernel_params.offset + optixGetLaunchIndex().x; + kernel_displace_evaluate(nullptr, input, output, 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; + const int global_index = kernel_params.offset + optixGetLaunchIndex().x; + kernel_background_evaluate(nullptr, input, output, global_index); +} + +extern "C" __global__ void __raygen__kernel_optix_shader_eval_curve_shadow_transparency() +{ + KernelShaderEvalInput *const input = (KernelShaderEvalInput *)kernel_params.path_index_array; + float *const output = kernel_params.render_buffer; + const int global_index = kernel_params.offset + optixGetLaunchIndex().x; + kernel_curve_shadow_transparency_evaluate(nullptr, input, output, global_index); +} diff --git a/intern/cycles/kernel/integrator/displacement_shader.h b/intern/cycles/kernel/integrator/displacement_shader.h index 839dfe244ac..a6e9d674396 100644 --- a/intern/cycles/kernel/integrator/displacement_shader.h +++ b/intern/cycles/kernel/integrator/displacement_shader.h @@ -24,8 +24,8 @@ ccl_device void displacement_shader_eval(KernelGlobals kg, /* this will modify sd->P */ #ifdef __OSL__ - if (kg->osl) { - OSLShader::eval_displacement(kg, state, sd); + if (kernel_data.kernel_features & KERNEL_FEATURE_OSL) { + osl_eval_nodes<SHADER_TYPE_DISPLACEMENT>(kg, state, sd, 0); } else #endif diff --git a/intern/cycles/kernel/integrator/surface_shader.h b/intern/cycles/kernel/integrator/surface_shader.h index 6c0097b11bd..5e47a34f77e 100644 --- a/intern/cycles/kernel/integrator/surface_shader.h +++ b/intern/cycles/kernel/integrator/surface_shader.h @@ -827,13 +827,8 @@ ccl_device void surface_shader_eval(KernelGlobals kg, sd->num_closure_left = max_closures; #ifdef __OSL__ - if (kg->osl) { - if (sd->object == OBJECT_NONE && sd->lamp == LAMP_NONE) { - OSLShader::eval_background(kg, state, sd, path_flag); - } - else { - OSLShader::eval_surface(kg, state, sd, path_flag); - } + if (kernel_data.kernel_features & KERNEL_FEATURE_OSL) { + osl_eval_nodes<SHADER_TYPE_SURFACE>(kg, state, sd, path_flag); } else #endif diff --git a/intern/cycles/kernel/integrator/volume_shader.h b/intern/cycles/kernel/integrator/volume_shader.h index 0ff968723a1..f9050647c6d 100644 --- a/intern/cycles/kernel/integrator/volume_shader.h +++ b/intern/cycles/kernel/integrator/volume_shader.h @@ -493,8 +493,8 @@ ccl_device_inline void volume_shader_eval(KernelGlobals kg, /* evaluate shader */ # ifdef __OSL__ - if (kg->osl) { - OSLShader::eval_volume(kg, state, sd, path_flag); + if (kernel_data.kernel_features & KERNEL_FEATURE_OSL) { + osl_eval_nodes<SHADER_TYPE_VOLUME>(kg, state, sd, path_flag); } else # endif diff --git a/intern/cycles/kernel/osl/closures.cpp b/intern/cycles/kernel/osl/closures.cpp index d56e0551a91..6800c765345 100644 --- a/intern/cycles/kernel/osl/closures.cpp +++ b/intern/cycles/kernel/osl/closures.cpp @@ -25,13 +25,18 @@ #include "kernel/osl/osl.h" -#include "kernel/osl/closures_setup.h" - #define TO_VEC3(v) OSL::Vec3(v.x, v.y, v.z) #define TO_FLOAT3(v) make_float3(v[0], v[1], v[2]) CCL_NAMESPACE_BEGIN +static_assert(sizeof(OSLClosure) == sizeof(OSL::ClosureColor) && + sizeof(OSLClosureAdd) == sizeof(OSL::ClosureAdd) && + sizeof(OSLClosureMul) == sizeof(OSL::ClosureMul) && + sizeof(OSLClosureComponent) == sizeof(OSL::ClosureComponent)); +static_assert(sizeof(ShaderGlobals) == sizeof(OSL::ShaderGlobals) && + offsetof(ShaderGlobals, Ci) == offsetof(OSL::ShaderGlobals, Ci)); + /* Registration */ #define OSL_CLOSURE_STRUCT_BEGIN(Upper, lower) \ @@ -60,53 +65,18 @@ void OSLRenderServices::register_closures(OSL::ShadingSystem *ss) #include "closures_template.h" } -/* Globals */ +/* Surface & Background */ -static void shaderdata_to_shaderglobals(const KernelGlobalsCPU *kg, - ShaderData *sd, - const void *state, - uint32_t path_flag, - OSLThreadData *tdata) +template<> +void osl_eval_nodes<SHADER_TYPE_SURFACE>(const KernelGlobalsCPU *kg, + const void *state, + ShaderData *sd, + uint32_t path_flag) { - OSL::ShaderGlobals *globals = &tdata->globals; - - const differential3 dP = differential_from_compact(sd->Ng, sd->dP); - const differential3 dI = differential_from_compact(sd->I, sd->dI); - - /* copy from shader data to shader globals */ - globals->P = TO_VEC3(sd->P); - globals->dPdx = TO_VEC3(dP.dx); - globals->dPdy = TO_VEC3(dP.dy); - globals->I = TO_VEC3(sd->I); - globals->dIdx = TO_VEC3(dI.dx); - globals->dIdy = TO_VEC3(dI.dy); - globals->N = TO_VEC3(sd->N); - globals->Ng = TO_VEC3(sd->Ng); - globals->u = sd->u; - globals->dudx = sd->du.dx; - globals->dudy = sd->du.dy; - globals->v = sd->v; - globals->dvdx = sd->dv.dx; - globals->dvdy = sd->dv.dy; - globals->dPdu = TO_VEC3(sd->dPdu); - globals->dPdv = TO_VEC3(sd->dPdv); - globals->surfacearea = 1.0f; - globals->time = sd->time; - - /* booleans */ - globals->raytype = path_flag; - globals->flipHandedness = 0; - globals->backfacing = (sd->flag & SD_BACKFACING); - - /* shader data to be used in services callbacks */ - globals->renderstate = sd; - - /* hacky, we leave it to services to fetch actual object matrix */ - globals->shader2common = sd; - globals->object2common = sd; - - /* must be set to NULL before execute */ - globals->Ci = NULL; + /* setup shader globals from shader data */ + OSLThreadData *tdata = kg->osl_tdata; + shaderdata_to_shaderglobals( + kg, sd, path_flag, reinterpret_cast<ShaderGlobals *>(&tdata->globals)); /* clear trace data */ tdata->tracedata.init = false; @@ -121,53 +91,6 @@ static void shaderdata_to_shaderglobals(const KernelGlobalsCPU *kg, sd->osl_path_state = (const IntegratorStateCPU *)state; sd->osl_shadow_path_state = nullptr; } -} - -static void flatten_closure_tree(const KernelGlobalsCPU *kg, - ShaderData *sd, - uint32_t path_flag, - const OSL::ClosureColor *closure, - float3 weight = make_float3(1.0f, 1.0f, 1.0f)) -{ - /* OSL gives us a closure tree, we flatten it into arrays per - * closure type, for evaluation, sampling, etc later on. */ - - switch (closure->id) { - case OSL::ClosureColor::MUL: { - OSL::ClosureMul *mul = (OSL::ClosureMul *)closure; - flatten_closure_tree(kg, sd, path_flag, mul->closure, TO_FLOAT3(mul->weight) * weight); - break; - } - case OSL::ClosureColor::ADD: { - OSL::ClosureAdd *add = (OSL::ClosureAdd *)closure; - flatten_closure_tree(kg, sd, path_flag, add->closureA, weight); - flatten_closure_tree(kg, sd, path_flag, add->closureB, weight); - break; - } -#define OSL_CLOSURE_STRUCT_BEGIN(Upper, lower) \ - case OSL_CLOSURE_##Upper##_ID: { \ - const OSL::ClosureComponent *comp = reinterpret_cast<const OSL::ClosureComponent *>(closure); \ - weight *= TO_FLOAT3(comp->w); \ - osl_closure_##lower##_setup( \ - kg, sd, path_flag, weight, reinterpret_cast<const Upper##Closure *>(comp + 1)); \ - break; \ - } -#include "closures_template.h" - default: - break; - } -} - -/* Surface */ - -void OSLShader::eval_surface(const KernelGlobalsCPU *kg, - const void *state, - ShaderData *sd, - uint32_t path_flag) -{ - /* setup shader globals from shader data */ - OSLThreadData *tdata = kg->osl_tdata; - shaderdata_to_shaderglobals(kg, sd, state, path_flag, tdata); /* execute shader for this point */ OSL::ShadingSystem *ss = (OSL::ShadingSystem *)kg->osl_ss; @@ -175,101 +98,99 @@ void OSLShader::eval_surface(const KernelGlobalsCPU *kg, OSL::ShadingContext *octx = tdata->context; int shader = sd->shader & SHADER_MASK; - /* automatic bump shader */ - if (kg->osl->bump_state[shader]) { - /* save state */ - const float3 P = sd->P; - const float dP = sd->dP; - const OSL::Vec3 dPdx = globals->dPdx; - const OSL::Vec3 dPdy = globals->dPdy; - - /* set state as if undisplaced */ - if (sd->flag & SD_HAS_DISPLACEMENT) { - float data[9]; - bool found = kg->osl->services->get_attribute(sd, - true, - OSLRenderServices::u_empty, - TypeDesc::TypeVector, - OSLRenderServices::u_geom_undisplaced, - data); - (void)found; - assert(found); - - differential3 tmp_dP; - memcpy(&sd->P, data, sizeof(float) * 3); - memcpy(&tmp_dP.dx, data + 3, sizeof(float) * 3); - memcpy(&tmp_dP.dy, data + 6, sizeof(float) * 3); - - object_position_transform(kg, sd, &sd->P); - object_dir_transform(kg, sd, &tmp_dP.dx); - object_dir_transform(kg, sd, &tmp_dP.dy); - - sd->dP = differential_make_compact(tmp_dP); - - globals->P = TO_VEC3(sd->P); - globals->dPdx = TO_VEC3(tmp_dP.dx); - globals->dPdy = TO_VEC3(tmp_dP.dy); + if (sd->object == OBJECT_NONE && sd->lamp == LAMP_NONE) { + /* background */ + if (kg->osl->background_state) { + ss->execute(octx, *(kg->osl->background_state), *globals); } - - /* execute bump shader */ - ss->execute(octx, *(kg->osl->bump_state[shader]), *globals); - - /* reset state */ - sd->P = P; - sd->dP = dP; - - globals->P = TO_VEC3(P); - globals->dPdx = TO_VEC3(dPdx); - globals->dPdy = TO_VEC3(dPdy); } + else { + /* automatic bump shader */ + if (kg->osl->bump_state[shader]) { + /* save state */ + const float3 P = sd->P; + const float dP = sd->dP; + const OSL::Vec3 dPdx = globals->dPdx; + const OSL::Vec3 dPdy = globals->dPdy; + + /* set state as if undisplaced */ + if (sd->flag & SD_HAS_DISPLACEMENT) { + float data[9]; + bool found = kg->osl->services->get_attribute(sd, + true, + OSLRenderServices::u_empty, + TypeDesc::TypeVector, + OSLRenderServices::u_geom_undisplaced, + data); + (void)found; + assert(found); + + differential3 tmp_dP; + memcpy(&sd->P, data, sizeof(float) * 3); + memcpy(&tmp_dP.dx, data + 3, sizeof(float) * 3); + memcpy(&tmp_dP.dy, data + 6, sizeof(float) * 3); + + object_position_transform(kg, sd, &sd->P); + object_dir_transform(kg, sd, &tmp_dP.dx); + object_dir_transform(kg, sd, &tmp_dP.dy); + + sd->dP = differential_make_compact(tmp_dP); + + globals->P = TO_VEC3(sd->P); + globals->dPdx = TO_VEC3(tmp_dP.dx); + globals->dPdy = TO_VEC3(tmp_dP.dy); + } + + /* execute bump shader */ + ss->execute(octx, *(kg->osl->bump_state[shader]), *globals); + + /* reset state */ + sd->P = P; + sd->dP = dP; + + globals->P = TO_VEC3(P); + globals->dPdx = TO_VEC3(dPdx); + globals->dPdy = TO_VEC3(dPdy); + } - /* surface shader */ - if (kg->osl->surface_state[shader]) { - ss->execute(octx, *(kg->osl->surface_state[shader]), *globals); + /* surface shader */ + if (kg->osl->surface_state[shader]) { + ss->execute(octx, *(kg->osl->surface_state[shader]), *globals); + } } /* flatten closure tree */ if (globals->Ci) { - flatten_closure_tree(kg, sd, path_flag, globals->Ci); + flatten_closure_tree(kg, sd, path_flag, reinterpret_cast<OSLClosure *>(globals->Ci)); } } -/* Background */ +/* Volume */ -void OSLShader::eval_background(const KernelGlobalsCPU *kg, - const void *state, - ShaderData *sd, - uint32_t path_flag) +template<> +void osl_eval_nodes<SHADER_TYPE_VOLUME>(const KernelGlobalsCPU *kg, + const void *state, + ShaderData *sd, + uint32_t path_flag) { /* setup shader globals from shader data */ OSLThreadData *tdata = kg->osl_tdata; - shaderdata_to_shaderglobals(kg, sd, state, path_flag, tdata); + shaderdata_to_shaderglobals( + kg, sd, path_flag, reinterpret_cast<ShaderGlobals *>(&tdata->globals)); - /* execute shader for this point */ - OSL::ShadingSystem *ss = (OSL::ShadingSystem *)kg->osl_ss; - OSL::ShaderGlobals *globals = &tdata->globals; - OSL::ShadingContext *octx = tdata->context; + /* clear trace data */ + tdata->tracedata.init = false; - if (kg->osl->background_state) { - ss->execute(octx, *(kg->osl->background_state), *globals); + /* Used by render-services. */ + sd->osl_globals = kg; + if (path_flag & PATH_RAY_SHADOW) { + sd->osl_path_state = nullptr; + sd->osl_shadow_path_state = (const IntegratorShadowStateCPU *)state; } - - /* return background color immediately */ - if (globals->Ci) { - flatten_closure_tree(kg, sd, path_flag, globals->Ci); + else { + sd->osl_path_state = (const IntegratorStateCPU *)state; + sd->osl_shadow_path_state = nullptr; } -} - -/* Volume */ - -void OSLShader::eval_volume(const KernelGlobalsCPU *kg, - const void *state, - ShaderData *sd, - uint32_t path_flag) -{ - /* setup shader globals from shader data */ - OSLThreadData *tdata = kg->osl_tdata; - shaderdata_to_shaderglobals(kg, sd, state, path_flag, tdata); /* execute shader */ OSL::ShadingSystem *ss = (OSL::ShadingSystem *)kg->osl_ss; @@ -283,17 +204,30 @@ void OSLShader::eval_volume(const KernelGlobalsCPU *kg, /* flatten closure tree */ if (globals->Ci) { - flatten_closure_tree(kg, sd, path_flag, globals->Ci); + flatten_closure_tree(kg, sd, path_flag, reinterpret_cast<OSLClosure *>(globals->Ci)); } } /* Displacement */ -void OSLShader::eval_displacement(const KernelGlobalsCPU *kg, const void *state, ShaderData *sd) +template<> +void osl_eval_nodes<SHADER_TYPE_DISPLACEMENT>(const KernelGlobalsCPU *kg, + const void *state, + ShaderData *sd, + uint32_t path_flag) { /* setup shader globals from shader data */ OSLThreadData *tdata = kg->osl_tdata; - shaderdata_to_shaderglobals(kg, sd, state, 0, tdata); + shaderdata_to_shaderglobals( + kg, sd, path_flag, reinterpret_cast<ShaderGlobals *>(&tdata->globals)); + + /* clear trace data */ + tdata->tracedata.init = false; + + /* Used by render-services. */ + sd->osl_globals = kg; + sd->osl_path_state = (const IntegratorStateCPU *)state; + sd->osl_shadow_path_state = nullptr; /* execute shader */ OSL::ShadingSystem *ss = (OSL::ShadingSystem *)kg->osl_ss; diff --git a/intern/cycles/kernel/osl/closures_setup.h b/intern/cycles/kernel/osl/closures_setup.h index 96c551b9951..ceaf56ccba6 100644 --- a/intern/cycles/kernel/osl/closures_setup.h +++ b/intern/cycles/kernel/osl/closures_setup.h @@ -40,12 +40,7 @@ CCL_NAMESPACE_BEGIN const char *label; #define OSL_CLOSURE_STRUCT_END(Upper, lower) \ } \ - ; \ - ccl_device void osl_closure_##lower##_setup(KernelGlobals kg, \ - ccl_private ShaderData *sd, \ - uint32_t path_flag, \ - float3 weight, \ - ccl_private Upper##Closure *closure); + ; #define OSL_CLOSURE_STRUCT_MEMBER(Upper, TYPE, type, name, key) type name; #define OSL_CLOSURE_STRUCT_ARRAY_MEMBER(Upper, TYPE, type, name, key, size) type name[size]; @@ -210,11 +205,9 @@ ccl_device void osl_closure_microfacet_setup(KernelGlobals kg, bsdf->ior = closure->ior; bsdf->T = closure->T; - static OSL::ustring u_ggx("ggx"); - static OSL::ustring u_default("default"); - /* GGX */ - if (closure->distribution == u_ggx || closure->distribution == u_default) { + if (closure->distribution == make_string("ggx", 11253504724482777663ull) || + closure->distribution == make_string("default", 4430693559278735917ull)) { if (!closure->refract) { if (closure->alpha_x == closure->alpha_y) { /* Isotropic */ @@ -1000,18 +993,14 @@ ccl_device void osl_closure_bssrdf_setup(KernelGlobals kg, float3 weight, ccl_private const BSSRDFClosure *closure) { - static ustring u_burley("burley"); - static ustring u_random_walk_fixed_radius("random_walk_fixed_radius"); - static ustring u_random_walk("random_walk"); - ClosureType type; - if (closure->method == u_burley) { + if (closure->method == make_string("burley", 186330084368958868ull)) { type = CLOSURE_BSSRDF_BURLEY_ID; } - else if (closure->method == u_random_walk_fixed_radius) { + else if (closure->method == make_string("random_walk_fixed_radius", 5695810351010063150ull)) { type = CLOSURE_BSSRDF_RANDOM_WALK_FIXED_RADIUS_ID; } - else if (closure->method == u_random_walk) { + else if (closure->method == make_string("random_walk", 11360609267673527222ull)) { type = CLOSURE_BSSRDF_RANDOM_WALK_ID; } else { diff --git a/intern/cycles/kernel/osl/closures_template.h b/intern/cycles/kernel/osl/closures_template.h index c808b275966..b9e9b52dcf8 100644 --- a/intern/cycles/kernel/osl/closures_template.h +++ b/intern/cycles/kernel/osl/closures_template.h @@ -40,7 +40,7 @@ OSL_CLOSURE_STRUCT_BEGIN(Transparent, transparent) OSL_CLOSURE_STRUCT_END(Transparent, transparent) OSL_CLOSURE_STRUCT_BEGIN(Microfacet, microfacet) - OSL_CLOSURE_STRUCT_MEMBER(Microfacet, STRING, ustring, distribution, NULL) + OSL_CLOSURE_STRUCT_MEMBER(Microfacet, STRING, DeviceString, distribution, NULL) OSL_CLOSURE_STRUCT_MEMBER(Microfacet, VECTOR, packed_float3, N, NULL) OSL_CLOSURE_STRUCT_MEMBER(Microfacet, VECTOR, packed_float3, T, NULL) OSL_CLOSURE_STRUCT_MEMBER(Microfacet, FLOAT, float, alpha_x, NULL) @@ -210,7 +210,7 @@ OSL_CLOSURE_STRUCT_BEGIN(PhongRamp, phong_ramp) OSL_CLOSURE_STRUCT_END(PhongRamp, phong_ramp) OSL_CLOSURE_STRUCT_BEGIN(BSSRDF, bssrdf) - OSL_CLOSURE_STRUCT_MEMBER(BSSRDF, STRING, ustring, method, NULL) + OSL_CLOSURE_STRUCT_MEMBER(BSSRDF, STRING, DeviceString, method, NULL) OSL_CLOSURE_STRUCT_MEMBER(BSSRDF, VECTOR, packed_float3, N, NULL) OSL_CLOSURE_STRUCT_MEMBER(BSSRDF, VECTOR, packed_float3, radius, NULL) OSL_CLOSURE_STRUCT_MEMBER(BSSRDF, VECTOR, packed_float3, albedo, NULL) diff --git a/intern/cycles/kernel/osl/osl.h b/intern/cycles/kernel/osl/osl.h index bef23f3eea1..cc5c81ad027 100644 --- a/intern/cycles/kernel/osl/osl.h +++ b/intern/cycles/kernel/osl/osl.h @@ -1,38 +1,171 @@ -/* SPDX-License-Identifier: Apache-2.0 - * Copyright 2011-2022 Blender Foundation */ +/* SPDX-License-Identifier: BSD-3-Clause + * + * Adapted from Open Shading Language + * Copyright (c) 2009-2010 Sony Pictures Imageworks Inc., et al. + * All Rights Reserved. + * + * Modifications Copyright 2011-2022 Blender Foundation. */ #pragma once /* OSL Shader Engine * - * Holds all variables to execute and use OSL shaders from the kernel. These - * are initialized externally by OSLShaderManager before rendering starts. - * - * Before/after a thread starts rendering, thread_init/thread_free must be - * called, which will store any per thread OSL state in thread local storage. - * This means no thread state must be passed along in the kernel itself. + * Holds all variables to execute and use OSL shaders from the kernel. */ #include "kernel/osl/types.h" +#include "kernel/osl/closures_setup.h" + CCL_NAMESPACE_BEGIN -class OSLShader { - public: - /* eval */ - static void eval_surface(const KernelGlobalsCPU *kg, - const void *state, - ShaderData *sd, - uint32_t path_flag); - static void eval_background(const KernelGlobalsCPU *kg, - const void *state, - ShaderData *sd, - uint32_t path_flag); - static void eval_volume(const KernelGlobalsCPU *kg, - const void *state, - ShaderData *sd, - uint32_t path_flag); - static void eval_displacement(const KernelGlobalsCPU *kg, const void *state, ShaderData *sd); -}; +ccl_device_inline void shaderdata_to_shaderglobals(KernelGlobals kg, + ccl_private ShaderData *sd, + uint32_t path_flag, + ccl_private ShaderGlobals *globals) +{ + const differential3 dP = differential_from_compact(sd->Ng, sd->dP); + const differential3 dI = differential_from_compact(sd->I, sd->dI); + + /* copy from shader data to shader globals */ + globals->P = sd->P; + globals->dPdx = dP.dx; + globals->dPdy = dP.dy; + globals->I = sd->I; + globals->dIdx = dI.dx; + globals->dIdy = dI.dy; + globals->N = sd->N; + globals->Ng = sd->Ng; + globals->u = sd->u; + globals->dudx = sd->du.dx; + globals->dudy = sd->du.dy; + globals->v = sd->v; + globals->dvdx = sd->dv.dx; + globals->dvdy = sd->dv.dy; + globals->dPdu = sd->dPdu; + globals->dPdv = sd->dPdv; + globals->time = sd->time; + globals->dtime = 1.0f; + globals->surfacearea = 1.0f; + globals->raytype = path_flag; + globals->flipHandedness = 0; + globals->backfacing = (sd->flag & SD_BACKFACING); + + /* shader data to be used in services callbacks */ + globals->renderstate = sd; + + /* hacky, we leave it to services to fetch actual object matrix */ + globals->shader2common = sd; + globals->object2common = sd; + + /* must be set to NULL before execute */ + globals->Ci = nullptr; +} + +ccl_device void flatten_closure_tree(KernelGlobals kg, + ccl_private ShaderData *sd, + uint32_t path_flag, + ccl_private const OSLClosure *closure) +{ + int stack_size = 0; + float3 weight = one_float3(); + float3 weight_stack[16]; + ccl_private const OSLClosure *closure_stack[16]; + + while (closure) { + switch (closure->id) { + case OSL_CLOSURE_MUL_ID: { + ccl_private const OSLClosureMul *mul = static_cast<ccl_private const OSLClosureMul *>( + closure); + weight *= mul->weight; + closure = mul->closure; + continue; + } + case OSL_CLOSURE_ADD_ID: { + if (stack_size >= 16) { + kernel_assert(!"Exhausted OSL closure stack"); + break; + } + ccl_private const OSLClosureAdd *add = static_cast<ccl_private const OSLClosureAdd *>( + closure); + closure = add->closureA; + weight_stack[stack_size] = weight; + closure_stack[stack_size++] = add->closureB; + continue; + } +#define OSL_CLOSURE_STRUCT_BEGIN(Upper, lower) \ + case OSL_CLOSURE_##Upper##_ID: { \ + ccl_private const OSLClosureComponent *comp = \ + static_cast<ccl_private const OSLClosureComponent *>(closure); \ + osl_closure_##lower##_setup(kg, \ + sd, \ + path_flag, \ + weight * comp->weight, \ + reinterpret_cast<ccl_private const Upper##Closure *>(comp + 1)); \ + break; \ + } +#include "closures_template.h" + default: + break; + } + + if (stack_size > 0) { + weight = weight_stack[--stack_size]; + closure = closure_stack[stack_size]; + } + else { + closure = nullptr; + } + } +} + +#ifndef __KERNEL_GPU__ + +template<ShaderType type> +void osl_eval_nodes(const KernelGlobalsCPU *kg, + const void *state, + ShaderData *sd, + uint32_t path_flag); + +#else + +template<ShaderType type, typename ConstIntegratorGenericState> +ccl_device_inline void osl_eval_nodes(KernelGlobals kg, + ConstIntegratorGenericState state, + ccl_private ShaderData *sd, + uint32_t path_flag) +{ + ShaderGlobals globals; + shaderdata_to_shaderglobals(kg, sd, path_flag, &globals); + + const int shader = sd->shader & SHADER_MASK; + +# ifdef __KERNEL_OPTIX__ + uint8_t group_data[2048]; + uint8_t closure_pool[1024]; + sd->osl_closure_pool = closure_pool; + + unsigned int optix_dc_index = 2 /* NUM_CALLABLE_PROGRAM_GROUPS */ + + (shader + type * kernel_data.max_shaders) * 2; + optixDirectCall<void>(optix_dc_index + 0, + /* shaderglobals_ptr = */ &globals, + /* groupdata_ptr = */ (void *)group_data, + /* userdata_base_ptr = */ (void *)nullptr, + /* output_base_ptr = */ (void *)nullptr, + /* shadeindex = */ 0); + optixDirectCall<void>(optix_dc_index + 1, + /* shaderglobals_ptr = */ &globals, + /* groupdata_ptr = */ (void *)group_data, + /* userdata_base_ptr = */ (void *)nullptr, + /* output_base_ptr = */ (void *)nullptr, + /* shadeindex = */ 0); +# endif + + if (globals.Ci) { + flatten_closure_tree(kg, sd, path_flag, globals.Ci); + } +} + +#endif CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/osl/services.cpp b/intern/cycles/kernel/osl/services.cpp index b744422ee78..454b75ea4d9 100644 --- a/intern/cycles/kernel/osl/services.cpp +++ b/intern/cycles/kernel/osl/services.cpp @@ -119,8 +119,8 @@ ustring OSLRenderServices::u_u("u"); ustring OSLRenderServices::u_v("v"); ustring OSLRenderServices::u_empty; -OSLRenderServices::OSLRenderServices(OSL::TextureSystem *texture_system) - : OSL::RendererServices(texture_system) +OSLRenderServices::OSLRenderServices(OSL::TextureSystem *texture_system, int device_type) + : OSL::RendererServices(texture_system), device_type_(device_type) { } @@ -131,6 +131,17 @@ OSLRenderServices::~OSLRenderServices() } } +int OSLRenderServices::supports(string_view feature) const +{ +#ifdef WITH_OPTIX + if (feature == "OptiX") { + return device_type_ == DEVICE_OPTIX; + } +#endif + + return false; +} + bool OSLRenderServices::get_matrix(OSL::ShaderGlobals *sg, OSL::Matrix44 &result, OSL::TransformationPtr xform, @@ -1139,29 +1150,39 @@ TextureSystem::TextureHandle *OSLRenderServices::get_texture_handle(ustring file { OSLTextureHandleMap::iterator it = textures.find(filename); - /* For non-OIIO textures, just return a pointer to our own OSLTextureHandle. */ - if (it != textures.end()) { - if (it->second->type != OSLTextureHandle::OIIO) { - return (TextureSystem::TextureHandle *)it->second.get(); + if (device_type_ == DEVICE_CPU) { + /* For non-OIIO textures, just return a pointer to our own OSLTextureHandle. */ + if (it != textures.end()) { + if (it->second->type != OSLTextureHandle::OIIO) { + return (TextureSystem::TextureHandle *)it->second.get(); + } } - } - /* Get handle from OpenImageIO. */ - OSL::TextureSystem *ts = m_texturesys; - TextureSystem::TextureHandle *handle = ts->get_texture_handle(filename); - if (handle == NULL) { - return NULL; - } + /* Get handle from OpenImageIO. */ + OSL::TextureSystem *ts = m_texturesys; + TextureSystem::TextureHandle *handle = ts->get_texture_handle(filename); + if (handle == NULL) { + return NULL; + } + + /* Insert new OSLTextureHandle if needed. */ + if (it == textures.end()) { + textures.insert(filename, new OSLTextureHandle(OSLTextureHandle::OIIO)); + it = textures.find(filename); + } - /* Insert new OSLTextureHandle if needed. */ - if (it == textures.end()) { - textures.insert(filename, new OSLTextureHandle(OSLTextureHandle::OIIO)); - it = textures.find(filename); + /* Assign OIIO texture handle and return. */ + it->second->oiio_handle = handle; + return (TextureSystem::TextureHandle *)it->second.get(); } + else { + if (it != textures.end() && it->second->type == OSLTextureHandle::SVM && it->second->svm_slots[0].w == -1) { + return reinterpret_cast<TextureSystem::TextureHandle *>( + static_cast<uintptr_t>(it->second->svm_slots[0].y + 1)); + } - /* Assign OIIO texture handle and return. */ - it->second->oiio_handle = handle; - return (TextureSystem::TextureHandle *)it->second.get(); + return NULL; + } } bool OSLRenderServices::good(TextureSystem::TextureHandle *texture_handle) diff --git a/intern/cycles/kernel/osl/services.h b/intern/cycles/kernel/osl/services.h index 334b6682e34..9d875ae8e94 100644 --- a/intern/cycles/kernel/osl/services.h +++ b/intern/cycles/kernel/osl/services.h @@ -22,11 +22,8 @@ class PtexCache; CCL_NAMESPACE_BEGIN -class Object; class Scene; -class Shader; struct ShaderData; -struct float3; struct KernelGlobalsCPU; /* OSL Texture Handle @@ -73,11 +70,13 @@ typedef OIIO::unordered_map_concurrent<ustring, OSLTextureHandleRef, ustringHash class OSLRenderServices : public OSL::RendererServices { public: - OSLRenderServices(OSL::TextureSystem *texture_system); + OSLRenderServices(OSL::TextureSystem *texture_system, int device_type); ~OSLRenderServices(); static void register_closures(OSL::ShadingSystem *ss); + int supports(string_view feature) const override; + bool get_matrix(OSL::ShaderGlobals *sg, OSL::Matrix44 &result, OSL::TransformationPtr xform, @@ -324,6 +323,9 @@ class OSLRenderServices : public OSL::RendererServices { * and is required because texture handles are cached as part of the shared * shading system. */ OSLTextureHandleMap textures; + + private: + int device_type_; }; CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/osl/services_gpu.h b/intern/cycles/kernel/osl/services_gpu.h new file mode 100644 index 00000000000..e6e19b8c484 --- /dev/null +++ b/intern/cycles/kernel/osl/services_gpu.h @@ -0,0 +1,2149 @@ +/* SPDX-License-Identifier: BSD-3-Clause + * + * Adapted from Open Shading Language + * Copyright (c) 2009-2010 Sony Pictures Imageworks Inc., et al. + * All Rights Reserved. + * + * Modifications Copyright 2011-2022 Blender Foundation. */ + +#include "kernel/tables.h" +#include "kernel/util/differential.h" + +#include "kernel/osl/osl.h" + +namespace DeviceStrings { + +/* "" */ +ccl_device_constant DeviceString _emptystring_ = {0ull}; +/* "NDC" */ +ccl_device_constant DeviceString u_ndc = {5148305047403260775ull}; +/* "screen" */ +ccl_device_constant DeviceString u_screen = {14159088609039777114ull}; +/* "camera" */ +ccl_device_constant DeviceString u_camera = {2159505832145726196ull}; +/* "raster" */ +ccl_device_constant DeviceString u_raster = {7759263238610201778ull}; +/* "world" */ +ccl_device_constant DeviceString u_world = {16436542438370751598ull}; +/* "common" */ +ccl_device_constant DeviceString u_common = {14645198576927606093ull}; +/* "hsv" */ +ccl_device_constant DeviceString u_hsv = {2177035556331879497ull}; +/* "hsl" */ +ccl_device_constant DeviceString u_hsl = {7749766809258288148ull}; +/* "XYZ" */ +ccl_device_constant DeviceString u_xyz = {4957977063494975483ull}; +/* "xyY" */ +ccl_device_constant DeviceString u_xyy = {5138822319725660255ull}; +/* "sRGB" */ +ccl_device_constant DeviceString u_srgb = {15368599878474175032ull}; +/* "object:location" */ +ccl_device_constant DeviceString u_object_location = {7846190347358762897ull}; +/* "object:color" */ +ccl_device_constant DeviceString u_object_color = {12695623857059169556ull}; +/* "object:alpha" */ +ccl_device_constant DeviceString u_object_alpha = {11165053919428293151ull}; +/* "object:index" */ +ccl_device_constant DeviceString u_object_index = {6588325838217472556ull}; +/* "geom:dupli_generated" */ +ccl_device_constant DeviceString u_geom_dupli_generated = {6715607178003388908ull}; +/* "geom:dupli_uv" */ +ccl_device_constant DeviceString u_geom_dupli_uv = {1294253317490155849ull}; +/* "material:index" */ +ccl_device_constant DeviceString u_material_index = {741770758159634623ull}; +/* "object:random" */ +ccl_device_constant DeviceString u_object_random = {15789063994977955884ull}; +/* "particle:index" */ +ccl_device_constant DeviceString u_particle_index = {9489711748229903784ull}; +/* "particle:random" */ +ccl_device_constant DeviceString u_particle_random = {17993722202766855761ull}; +/* "particle:age" */ +ccl_device_constant DeviceString u_particle_age = {7380730644710951109ull}; +/* "particle:lifetime" */ +ccl_device_constant DeviceString u_particle_lifetime = {16576828923156200061ull}; +/* "particle:location" */ +ccl_device_constant DeviceString u_particle_location = {10309536211423573010ull}; +/* "particle:rotation" */ +ccl_device_constant DeviceString u_particle_rotation = {17858543768041168459ull}; +/* "particle:size" */ +ccl_device_constant DeviceString u_particle_size = {16461524249715420389ull}; +/* "particle:velocity" */ +ccl_device_constant DeviceString u_particle_velocity = {13199101248768308863ull}; +/* "particle:angular_velocity" */ +ccl_device_constant DeviceString u_particle_angular_velocity = {16327930120486517910ull}; +/* "geom:numpolyvertices" */ +ccl_device_constant DeviceString u_geom_numpolyvertices = {382043551489988826ull}; +/* "geom:trianglevertices" */ +ccl_device_constant DeviceString u_geom_trianglevertices = {17839267571524187074ull}; +/* "geom:polyvertices" */ +ccl_device_constant DeviceString u_geom_polyvertices = {1345577201967881769ull}; +/* "geom:name" */ +ccl_device_constant DeviceString u_geom_name = {13606338128269760050ull}; +/* "geom:undisplaced" */ +ccl_device_constant DeviceString u_geom_undisplaced = {12431586303019276305ull}; +/* "geom:is_smooth" */ +ccl_device_constant DeviceString u_is_smooth = {857544214094480123ull}; +/* "geom:is_curve" */ +ccl_device_constant DeviceString u_is_curve = {129742495633653138ull}; +/* "geom:curve_thickness" */ +ccl_device_constant DeviceString u_curve_thickness = {10605802038397633852ull}; +/* "geom:curve_length" */ +ccl_device_constant DeviceString u_curve_length = {11423459517663715453ull}; +/* "geom:curve_tangent_normal" */ +ccl_device_constant DeviceString u_curve_tangent_normal = {12301397394034985633ull}; +/* "geom:curve_random" */ +ccl_device_constant DeviceString u_curve_random = {15293085049960492358ull}; +/* "geom:is_point" */ +ccl_device_constant DeviceString u_is_point = {2511357849436175953ull}; +/* "geom:point_radius" */ +ccl_device_constant DeviceString u_point_radius = {9956381140398668479ull}; +/* "geom:point_position" */ +ccl_device_constant DeviceString u_point_position = {15684484280742966916ull}; +/* "geom:point_random" */ +ccl_device_constant DeviceString u_point_random = {5632627207092325544ull}; +/* "geom:normal_map_normal" */ +ccl_device_constant DeviceString u_normal_map_normal = {10718948685686827073}; +/* "path:ray_length" */ +ccl_device_constant DeviceString u_path_ray_length = {16391985802412544524ull}; +/* "path:ray_depth" */ +ccl_device_constant DeviceString u_path_ray_depth = {16643933224879500399ull}; +/* "path:diffuse_depth" */ +ccl_device_constant DeviceString u_path_diffuse_depth = {13191651286699118408ull}; +/* "path:glossy_depth" */ +ccl_device_constant DeviceString u_path_glossy_depth = {15717768399057252940ull}; +/* "path:transparent_depth" */ +ccl_device_constant DeviceString u_path_transparent_depth = {7821650266475578543ull}; +/* "path:transmission_depth" */ +ccl_device_constant DeviceString u_path_transmission_depth = {15113408892323917624ull}; + +} // namespace DeviceStrings + +/* Closure */ + +ccl_device_extern ccl_private OSLClosure *osl_mul_closure_color(ccl_private ShaderGlobals *sg, + ccl_private OSLClosure *a, + ccl_private const float3 *weight) +{ + if (*weight == zero_float3() || !a) { + return nullptr; + } + else if (*weight == one_float3()) { + return a; + } + + ccl_private ShaderData *const sd = static_cast<ccl_private ShaderData *>(sg->renderstate); + + ccl_private uint8_t *closure_pool = sd->osl_closure_pool; + /* Align pointer to closure struct requirement */ + closure_pool = reinterpret_cast<uint8_t *>( + (reinterpret_cast<size_t>(closure_pool) + alignof(OSLClosureMul) - 1) & + (-alignof(OSLClosureMul))); + sd->osl_closure_pool = closure_pool + sizeof(OSLClosureMul); + + ccl_private OSLClosureMul *const closure = reinterpret_cast<ccl_private OSLClosureMul *>( + closure_pool); + closure->id = OSL_CLOSURE_MUL_ID; + closure->weight = *weight; + closure->closure = a; + + return closure; +} + +ccl_device_extern ccl_private OSLClosure *osl_mul_closure_float(ccl_private ShaderGlobals *sg, + ccl_private OSLClosure *a, + float weight) +{ + if (weight == 0.0f || !a) { + return nullptr; + } + else if (weight == 1.0f) { + return a; + } + + ccl_private ShaderData *const sd = static_cast<ccl_private ShaderData *>(sg->renderstate); + + uint8_t *closure_pool = sd->osl_closure_pool; + /* Align pointer to closure struct requirement */ + closure_pool = reinterpret_cast<uint8_t *>( + (reinterpret_cast<size_t>(closure_pool) + alignof(OSLClosureMul) - 1) & + (-alignof(OSLClosureMul))); + sd->osl_closure_pool = closure_pool + sizeof(OSLClosureMul); + + ccl_private OSLClosureMul *const closure = reinterpret_cast<ccl_private OSLClosureMul *>( + closure_pool); + closure->id = OSL_CLOSURE_MUL_ID; + closure->weight = make_float3(weight, weight, weight); + closure->closure = a; + + return closure; +} + +ccl_device_extern ccl_private OSLClosure *osl_add_closure_closure(ccl_private ShaderGlobals *sg, + ccl_private OSLClosure *a, + ccl_private OSLClosure *b) +{ + if (!a) { + return b; + } + if (!b) { + return a; + } + + ccl_private ShaderData *const sd = static_cast<ccl_private ShaderData *>(sg->renderstate); + + ccl_private uint8_t *closure_pool = sd->osl_closure_pool; + /* Align pointer to closure struct requirement */ + closure_pool = reinterpret_cast<uint8_t *>( + (reinterpret_cast<size_t>(closure_pool) + alignof(OSLClosureAdd) - 1) & + (-alignof(OSLClosureAdd))); + sd->osl_closure_pool = closure_pool + sizeof(OSLClosureAdd); + + ccl_private OSLClosureAdd *const closure = reinterpret_cast<ccl_private OSLClosureAdd *>( + closure_pool); + closure->id = OSL_CLOSURE_ADD_ID; + closure->closureA = a; + closure->closureB = b; + + return closure; +} + +ccl_device_extern ccl_private OSLClosure *osl_allocate_closure_component( + ccl_private ShaderGlobals *sg, int id, int size) +{ + ccl_private ShaderData *const sd = static_cast<ccl_private ShaderData *>(sg->renderstate); + + ccl_private uint8_t *closure_pool = sd->osl_closure_pool; + /* Align pointer to closure struct requirement */ + closure_pool = reinterpret_cast<uint8_t *>( + (reinterpret_cast<size_t>(closure_pool) + alignof(OSLClosureComponent) - 1) & + (-alignof(OSLClosureComponent))); + sd->osl_closure_pool = closure_pool + sizeof(OSLClosureComponent) + size; + + ccl_private OSLClosureComponent *const closure = + reinterpret_cast<ccl_private OSLClosureComponent *>(closure_pool); + closure->id = static_cast<OSLClosureType>(id); + closure->weight = one_float3(); + + return closure; +} + +ccl_device_extern ccl_private OSLClosure *osl_allocate_weighted_closure_component( + ccl_private ShaderGlobals *sg, int id, int size, ccl_private const float3 *weight) +{ + ccl_private ShaderData *const sd = static_cast<ccl_private ShaderData *>(sg->renderstate); + + ccl_private uint8_t *closure_pool = sd->osl_closure_pool; + /* Align pointer to closure struct requirement */ + closure_pool = reinterpret_cast<uint8_t *>( + (reinterpret_cast<size_t>(closure_pool) + alignof(OSLClosureComponent) - 1) & + (-alignof(OSLClosureComponent))); + sd->osl_closure_pool = closure_pool + sizeof(OSLClosureComponent) + size; + + ccl_private OSLClosureComponent *const closure = + reinterpret_cast<ccl_private OSLClosureComponent *>(closure_pool); + closure->id = static_cast<OSLClosureType>(id); + closure->weight = *weight; + + return closure; +} + +/* Utilities */ + +#include "kernel/svm/math_util.h" +#include "kernel/util/color.h" + +ccl_device_extern void osl_error(ccl_private ShaderGlobals *sg, const char *format, void *args) +{ +} + +ccl_device_extern void osl_printf(ccl_private ShaderGlobals *sg, const char *format, void *args) +{ +} + +ccl_device_extern void osl_warning(ccl_private ShaderGlobals *sg, const char *format, void *args) +{ +} + +ccl_device_extern uint osl_range_check(int indexvalue, + int length, + DeviceString symname, + ccl_private ShaderGlobals *sg, + DeviceString sourcefile, + int sourceline, + DeviceString groupname, + int layer, + DeviceString layername, + DeviceString shadername) +{ + const int result = indexvalue < 0 ? 0 : indexvalue >= length ? length - 1 : indexvalue; +#if 0 + if (result != indexvalue) { + printf("Index [%d] out of range\n", indexvalue); + } +#endif + return result; +} + +ccl_device_extern uint osl_range_check_err(int indexvalue, + int length, + DeviceString symname, + ccl_private ShaderGlobals *sg, + DeviceString sourcefile, + int sourceline, + DeviceString groupname, + int layer, + DeviceString layername, + DeviceString shadername) +{ + return osl_range_check(indexvalue, + length, + symname, + sg, + sourcefile, + sourceline, + groupname, + layer, + layername, + shadername); +} + +/* Color Utilities */ + +ccl_device_extern void osl_blackbody_vf(ccl_private ShaderGlobals *sg, + ccl_private float3 *result, + float temperature) +{ + float3 color_rgb = rec709_to_rgb(nullptr, svm_math_blackbody_color_rec709(temperature)); + color_rgb = max(color_rgb, zero_float3()); + *result = color_rgb; +} + +#if 0 +ccl_device_extern void osl_wavelength_color_vf(ccl_private ShaderGlobals *sg, + ccl_private float3 *result, + float wavelength) +{ +} +#endif + +ccl_device_extern void osl_luminance_fv(ccl_private ShaderGlobals *sg, + ccl_private float *result, + ccl_private float3 *color) +{ + *result = linear_rgb_to_gray(nullptr, *color); +} + +ccl_device_extern void osl_luminance_dfdv(ccl_private ShaderGlobals *sg, + ccl_private float *result, + ccl_private float3 *color) +{ + for (int i = 0; i < 3; ++i) { + osl_luminance_fv(sg, result + i, color + i); + } +} + +ccl_device_extern void osl_prepend_color_from(ccl_private ShaderGlobals *sg, + ccl_private float3 *res, + DeviceString from) +{ + if (from == DeviceStrings::u_hsv) { + *res = hsv_to_rgb(*res); + } + else if (from == DeviceStrings::u_hsl) { + *res = hsl_to_rgb(*res); + } + else if (from == DeviceStrings::u_xyz) { + *res = xyz_to_rgb(nullptr, *res); + } + else if (from == DeviceStrings::u_xyy) { + *res = xyz_to_rgb(nullptr, xyY_to_xyz(res->x, res->y, res->z)); + } +} + +ccl_device_extern bool osl_transformc(ccl_private ShaderGlobals *sg, + ccl_private float3 *c_in, + int c_in_derivs, + ccl_private float3 *c_out, + int c_out_derivs, + DeviceString from, + DeviceString to) +{ + if (!c_out_derivs) { + c_in_derivs = false; + } + else if (!c_in_derivs) { + c_out[1] = zero_float3(); + c_out[2] = zero_float3(); + } + + float3 rgb; + + for (int i = 0; i < (c_in_derivs ? 3 : 1); ++i) { + if (from == DeviceStrings::u_hsv) { + rgb = hsv_to_rgb(c_in[i]); + } + else if (from == DeviceStrings::u_hsl) { + rgb = hsl_to_rgb(c_in[i]); + } + else if (from == DeviceStrings::u_xyz) { + rgb = xyz_to_rgb(nullptr, c_in[i]); + } + else if (from == DeviceStrings::u_xyy) { + rgb = xyz_to_rgb(nullptr, xyY_to_xyz(c_in[i].x, c_in[i].y, c_in[i].z)); + } + else if (from == DeviceStrings::u_srgb) { + rgb = color_srgb_to_linear_v3(c_in[i]); + } + else { + rgb = c_in[i]; + } + + if (to == DeviceStrings::u_hsv) { + c_out[i] = rgb_to_hsv(rgb); + } + else if (to == DeviceStrings::u_hsl) { + c_out[i] = rgb_to_hsl(rgb); + } +#if 0 + else if (to == DeviceStrings::u_xyz) { + c_out[i] = rgb_to_xyz(nullptr, rgb); + } + else if (to == DeviceStrings::u_xyy) { + c_out[i] = xyz_to_xyY(rgb_to_xyz(nullptr, rgb)); + } +#endif + else if (to == DeviceStrings::u_srgb) { + c_out[i] = color_linear_to_srgb_v3(rgb); + } + else { + c_out[i] = rgb; + } + } +} + +/* Matrix Utilities */ + +#include "util/transform.h" + +ccl_device_forceinline void copy_matrix(ccl_private float *res, const Transform &tfm) +{ + res[0] = tfm.x.x; + res[1] = tfm.y.x; + res[2] = tfm.z.x; + res[3] = 0.0f; + res[4] = tfm.x.y; + res[5] = tfm.y.y; + res[6] = tfm.z.y; + res[7] = 0.0f; + res[8] = tfm.x.z; + res[9] = tfm.y.z; + res[10] = tfm.z.z; + res[11] = 0.0f; + res[12] = tfm.x.w; + res[13] = tfm.y.w; + res[14] = tfm.z.w; + res[15] = 1.0f; +} +ccl_device_forceinline void copy_matrix(ccl_private float *res, const ProjectionTransform &tfm) +{ + res[0] = tfm.x.x; + res[1] = tfm.y.x; + res[2] = tfm.z.x; + res[3] = tfm.w.x; + res[4] = tfm.x.y; + res[5] = tfm.y.y; + res[6] = tfm.z.y; + res[7] = tfm.w.y; + res[8] = tfm.x.z; + res[9] = tfm.y.z; + res[10] = tfm.z.z; + res[11] = tfm.w.z; + res[12] = tfm.x.w; + res[13] = tfm.y.w; + res[14] = tfm.z.w; + res[15] = tfm.w.w; +} +ccl_device_forceinline void copy_identity_matrix(ccl_private float *res) +{ + res[0] = 1.0f; + res[1] = 0.0f; + res[2] = 0.0f; + res[3] = 0.0f; + res[4] = 0.0f; + res[5] = 1.0f; + res[6] = 0.0f; + res[7] = 0.0f; + res[8] = 0.0f; + res[9] = 0.0f; + res[10] = 1.0f; + res[11] = 0.0f; + res[12] = 0.0f; + res[13] = 0.0f; + res[14] = 0.0f; + res[15] = 1.0f; +} +ccl_device_forceinline Transform convert_transform(ccl_private const float *m) +{ + return make_transform( + m[0], m[4], m[8], m[12], m[1], m[5], m[9], m[13], m[2], m[6], m[10], m[14]); +} + +ccl_device_extern void osl_mul_mmm(ccl_private float *res, + ccl_private const float *a, + ccl_private const float *b) +{ + const Transform tfm_a = convert_transform(a); + const Transform tfm_b = convert_transform(b); + copy_matrix(res, tfm_a * tfm_b); +} + +ccl_device_extern void osl_mul_mmf(ccl_private float *res, ccl_private const float *a, float b) +{ + for (int i = 0; i < 16; ++i) { + res[i] = a[i] * b; + } +} + +ccl_device_extern void osl_div_mmm(ccl_private float *res, + ccl_private const float *a, + ccl_private const float *b) +{ + const Transform tfm_a = convert_transform(a); + const Transform tfm_b = convert_transform(b); + copy_matrix(res, tfm_a * transform_inverse(tfm_b)); +} + +ccl_device_extern void osl_div_mmf(ccl_private float *res, ccl_private const float *a, float b) +{ + for (int i = 0; i < 16; ++i) { + res[i] = a[i] / b; + } +} + +ccl_device_extern void osl_div_mfm(ccl_private float *res, float a, ccl_private const float *b) +{ + const Transform tfm_b = convert_transform(b); + copy_matrix(res, transform_inverse(tfm_b)); + for (int i = 0; i < 16; ++i) { + res[i] *= a; + } +} + +ccl_device_extern void osl_div_m_ff(ccl_private float *res, float a, float b) +{ + float f = (b == 0) ? 0.0f : (a / b); + res[0] = f; + res[1] = 0.0f; + res[2] = 0.0f; + res[3] = 0.0f; + res[4] = 0.0f; + res[5] = f; + res[6] = 0.0f; + res[7] = 0.0f; + res[8] = 0.0f; + res[9] = 0.0f; + res[10] = f; + res[11] = 0.0f; + res[12] = 0.0f; + res[13] = 0.0f; + res[14] = 0.0f; + res[15] = f; +} + +ccl_device_extern void osl_transform_vmv(ccl_private float3 *res, + ccl_private const float *m, + ccl_private const float3 *v) +{ + const Transform tfm_m = convert_transform(m); + *res = transform_point(&tfm_m, *v); +} + +ccl_device_extern void osl_transform_dvmdv(ccl_private float3 *res, + ccl_private const float *m, + ccl_private const float3 *v) +{ + for (int i = 0; i < 3; ++i) { + const Transform tfm_m = convert_transform(m + i * 16); + res[i] = transform_point(&tfm_m, v[i]); + } +} + +ccl_device_extern void osl_transformv_vmv(ccl_private float3 *res, + ccl_private const float *m, + ccl_private const float3 *v) +{ + const Transform tfm_m = convert_transform(m); + *res = transform_direction(&tfm_m, *v); +} + +ccl_device_extern void osl_transformv_dvmdv(ccl_private float3 *res, + ccl_private const float *m, + ccl_private const float3 *v) +{ + for (int i = 0; i < 3; ++i) { + const Transform tfm_m = convert_transform(m + i * 16); + res[i] = transform_direction(&tfm_m, v[i]); + } +} + +ccl_device_extern void osl_transformn_vmv(ccl_private float3 *res, + ccl_private const float *m, + ccl_private const float3 *v) +{ + const Transform tfm_m = convert_transform(m); + *res = transform_direction(&tfm_m, *v); +} + +ccl_device_extern void osl_transformn_dvmdv(ccl_private float3 *res, + ccl_private const float *m, + ccl_private const float3 *v) +{ + for (int i = 0; i < 3; ++i) { + const Transform tfm_m = convert_transform(m + i * 16); + res[i] = transform_direction(&tfm_m, v[i]); + } +} + +ccl_device_extern bool osl_get_matrix(ccl_private ShaderGlobals *sg, + ccl_private float *result, + DeviceString from) +{ + if (from == DeviceStrings::u_ndc) { + copy_matrix(result, kernel_data.cam.ndctoworld); + return true; + } + if (from == DeviceStrings::u_raster) { + copy_matrix(result, kernel_data.cam.rastertoworld); + return true; + } + if (from == DeviceStrings::u_screen) { + copy_matrix(result, kernel_data.cam.screentoworld); + return true; + } + if (from == DeviceStrings::u_camera) { + copy_matrix(result, kernel_data.cam.cameratoworld); + return true; + } + if (from == DeviceStrings::u_world) { + copy_identity_matrix(result); + return true; + } + + return false; +} + +ccl_device_extern bool osl_get_inverse_matrix(ccl_private ShaderGlobals *sg, + ccl_private float *res, + DeviceString to) +{ + if (to == DeviceStrings::u_ndc) { + copy_matrix(res, kernel_data.cam.worldtondc); + return true; + } + if (to == DeviceStrings::u_raster) { + copy_matrix(res, kernel_data.cam.worldtoraster); + return true; + } + if (to == DeviceStrings::u_screen) { + copy_matrix(res, kernel_data.cam.worldtoscreen); + return true; + } + if (to == DeviceStrings::u_camera) { + copy_matrix(res, kernel_data.cam.worldtocamera); + return true; + } + if (to == DeviceStrings::u_world) { + copy_identity_matrix(res); + return true; + } + + return false; +} + +ccl_device_extern bool osl_get_from_to_matrix(ccl_private ShaderGlobals *sg, + ccl_private float *res, + DeviceString from, + DeviceString to) +{ + float m_from[16], m_to[16]; + if (osl_get_matrix(sg, m_from, from) && osl_get_inverse_matrix(sg, m_to, to)) { + osl_mul_mmm(res, m_from, m_to); + return true; + } + + return false; +} + +ccl_device_extern void osl_prepend_matrix_from(ccl_private ShaderGlobals *sg, + ccl_private float *res, + DeviceString from) +{ + float m[16]; + if (osl_get_matrix(sg, m, from)) { + osl_mul_mmm(res, m, res); + } +} + +ccl_device_extern bool osl_transform_triple(ccl_private ShaderGlobals *sg, + ccl_private float3 *p_in, + int p_in_derivs, + ccl_private float3 *p_out, + int p_out_derivs, + DeviceString from, + DeviceString to, + int vectype) +{ + if (!p_out_derivs) { + p_in_derivs = false; + } + else if (!p_in_derivs) { + p_out[1] = zero_float3(); + p_out[2] = zero_float3(); + } + + bool res; + float m[16]; + + if (from == DeviceStrings::u_common) { + res = osl_get_inverse_matrix(sg, m, to); + } + else if (to == DeviceStrings::u_common) { + res = osl_get_matrix(sg, m, from); + } + else { + res = osl_get_from_to_matrix(sg, m, from, to); + } + + if (res) { + if (vectype == 2 /* TypeDesc::POINT */) { + if (p_in_derivs) + osl_transform_dvmdv(p_out, m, p_in); + else + osl_transform_vmv(p_out, m, p_in); + } + else if (vectype == 3 /* TypeDesc::VECTOR */) { + if (p_in_derivs) + osl_transformv_dvmdv(p_out, m, p_in); + else + osl_transformv_vmv(p_out, m, p_in); + } + else if (vectype == 4 /* TypeDesc::NORMAL */) { + if (p_in_derivs) + osl_transformn_dvmdv(p_out, m, p_in); + else + osl_transformn_vmv(p_out, m, p_in); + } + else { + res = false; + } + } + else { + p_out[0] = p_in[0]; + if (p_in_derivs) { + p_out[1] = p_in[1]; + p_out[2] = p_in[2]; + } + } + + return res; +} + +ccl_device_extern bool osl_transform_triple_nonlinear(ccl_private ShaderGlobals *sg, + ccl_private float3 *p_in, + int p_in_derivs, + ccl_private float3 *p_out, + int p_out_derivs, + DeviceString from, + DeviceString to, + int vectype) +{ + return osl_transform_triple(sg, p_in, p_in_derivs, p_out, p_out_derivs, from, to, vectype); +} + +ccl_device_extern void osl_transpose_mm(ccl_private float *res, ccl_private const float *m) +{ + copy_matrix(res, *reinterpret_cast<ccl_private const ProjectionTransform *>(m)); +} + +#if 0 +ccl_device_extern float osl_determinant_fm(ccl_private const float *m) +{ +} +#endif + +/* Attributes */ + +#include "kernel/geom/geom.h" + +typedef long long TypeDesc; + +ccl_device_inline bool set_attribute_float(ccl_private float fval[3], + TypeDesc type, + bool derivatives, + ccl_private void *val) +{ + const unsigned char type_basetype = type & 0xF; + const unsigned char type_aggregate = (type >> 8) & 0xF; + const int type_arraylen = type >> 32; + + if (type_basetype == 11 /* TypeDesc::FLOAT */) { + if ((type_aggregate == 2 /* TypeDesc::VEC2 */) || + (type_aggregate == 1 && type_arraylen == 2)) { + for (int i = 0; i < (derivatives ? 3 : 1); ++i) { + static_cast<ccl_private float *>(val)[i * 2 + 0] = fval[i]; + static_cast<ccl_private float *>(val)[i * 2 + 1] = fval[i]; + } + return true; + } + if ((type_aggregate == 3 /* TypeDesc::VEC3 */) || + (type_aggregate == 1 && type_arraylen == 3)) { + for (int i = 0; i < (derivatives ? 3 : 1); ++i) { + static_cast<ccl_private float *>(val)[i * 3 + 0] = fval[i]; + static_cast<ccl_private float *>(val)[i * 3 + 1] = fval[i]; + static_cast<ccl_private float *>(val)[i * 3 + 2] = fval[i]; + } + return true; + } + if ((type_aggregate == 4 /* TypeDesc::VEC4 */) || + (type_aggregate == 1 && type_arraylen == 4)) { + for (int i = 0; i < (derivatives ? 3 : 1); ++i) { + static_cast<ccl_private float *>(val)[i * 4 + 0] = fval[i]; + static_cast<ccl_private float *>(val)[i * 4 + 1] = fval[i]; + static_cast<ccl_private float *>(val)[i * 4 + 2] = fval[i]; + static_cast<ccl_private float *>(val)[i * 4 + 3] = 1.0f; + } + return true; + } + if ((type_aggregate == 1 /* TypeDesc::SCALAR */)) { + for (int i = 0; i < (derivatives ? 3 : 1); ++i) { + static_cast<ccl_private float *>(val)[i] = fval[i]; + } + return true; + } + } + + return false; +} +ccl_device_inline bool set_attribute_float(float f, + TypeDesc type, + bool derivatives, + ccl_private void *val) +{ + float fv[3]; + + fv[0] = f; + fv[1] = 0.0f; + fv[2] = 0.0f; + + return set_attribute_float(fv, type, derivatives, val); +} +ccl_device_inline bool set_attribute_float2(ccl_private float2 fval[3], + TypeDesc type, + bool derivatives, + ccl_private void *val) +{ + const unsigned char type_basetype = type & 0xF; + const unsigned char type_aggregate = (type >> 8) & 0xF; + const int type_arraylen = type >> 32; + + if (type_basetype == 11 /* TypeDesc::FLOAT */) { + if ((type_aggregate == 2 /* TypeDesc::VEC2 */) || + (type_aggregate == 1 && type_arraylen == 2)) { + for (int i = 0; i < (derivatives ? 3 : 1); ++i) { + static_cast<ccl_private float *>(val)[i * 2 + 0] = fval[i].x; + static_cast<ccl_private float *>(val)[i * 2 + 1] = fval[i].y; + } + return true; + } + if ((type_aggregate == 3 /* TypeDesc::VEC3 */) || + (type_aggregate == 1 && type_arraylen == 3)) { + for (int i = 0; i < (derivatives ? 3 : 1); ++i) { + static_cast<ccl_private float *>(val)[i * 3 + 0] = fval[i].x; + static_cast<ccl_private float *>(val)[i * 3 + 1] = fval[i].y; + static_cast<ccl_private float *>(val)[i * 3 + 2] = 0.0f; + } + return true; + } + if ((type_aggregate == 4 /* TypeDesc::VEC4 */) || + (type_aggregate == 1 && type_arraylen == 4)) { + for (int i = 0; i < (derivatives ? 3 : 1); ++i) { + static_cast<ccl_private float *>(val)[i * 4 + 0] = fval[i].x; + static_cast<ccl_private float *>(val)[i * 4 + 1] = fval[i].y; + static_cast<ccl_private float *>(val)[i * 4 + 2] = 0.0f; + static_cast<ccl_private float *>(val)[i * 4 + 3] = 1.0f; + } + return true; + } + if ((type_aggregate == 1 /* TypeDesc::SCALAR */)) { + for (int i = 0; i < (derivatives ? 3 : 1); ++i) { + static_cast<ccl_private float *>(val)[i] = fval[i].x; + } + return true; + } + } + + return false; +} +ccl_device_inline bool set_attribute_float3(ccl_private float3 fval[3], + TypeDesc type, + bool derivatives, + ccl_private void *val) +{ + const unsigned char type_basetype = type & 0xF; + const unsigned char type_aggregate = (type >> 8) & 0xF; + const int type_arraylen = type >> 32; + + if (type_basetype == 11 /* TypeDesc::FLOAT */) { + if ((type_aggregate == 3 /* TypeDesc::VEC3 */) || + (type_aggregate == 1 && type_arraylen == 3)) { + for (int i = 0; i < (derivatives ? 3 : 1); ++i) { + static_cast<ccl_private float *>(val)[i * 3 + 0] = fval[i].x; + static_cast<ccl_private float *>(val)[i * 3 + 1] = fval[i].y; + static_cast<ccl_private float *>(val)[i * 3 + 2] = fval[i].z; + } + return true; + } + if ((type_aggregate == 4 /* TypeDesc::VEC4 */) || + (type_aggregate == 1 && type_arraylen == 4)) { + for (int i = 0; i < (derivatives ? 3 : 1); ++i) { + static_cast<ccl_private float *>(val)[i * 4 + 0] = fval[i].x; + static_cast<ccl_private float *>(val)[i * 4 + 1] = fval[i].y; + static_cast<ccl_private float *>(val)[i * 4 + 2] = fval[i].z; + static_cast<ccl_private float *>(val)[i * 4 + 3] = 1.0f; + } + return true; + } + if ((type_aggregate == 1 /* TypeDesc::SCALAR */)) { + for (int i = 0; i < (derivatives ? 3 : 1); ++i) { + static_cast<ccl_private float *>(val)[i] = average(fval[i]); + } + return true; + } + } + + return false; +} +ccl_device_inline bool set_attribute_float3(float3 f, + TypeDesc type, + bool derivatives, + ccl_private void *val) +{ + float3 fv[3]; + + fv[0] = f; + fv[1] = make_float3(0.0f, 0.0f, 0.0f); + fv[2] = make_float3(0.0f, 0.0f, 0.0f); + + return set_attribute_float3(fv, type, derivatives, val); +} +ccl_device_inline bool set_attribute_float4(ccl_private float4 fval[3], + TypeDesc type, + bool derivatives, + ccl_private void *val) +{ + const unsigned char type_basetype = type & 0xF; + const unsigned char type_aggregate = (type >> 8) & 0xF; + const int type_arraylen = type >> 32; + + if (type_basetype == 11 /* TypeDesc::FLOAT */) { + if ((type_aggregate == 3 /* TypeDesc::VEC3 */) || + (type_aggregate == 1 && type_arraylen == 3)) { + for (int i = 0; i < (derivatives ? 3 : 1); ++i) { + static_cast<ccl_private float *>(val)[i * 3 + 0] = fval[i].x; + static_cast<ccl_private float *>(val)[i * 3 + 1] = fval[i].y; + static_cast<ccl_private float *>(val)[i * 3 + 2] = fval[i].z; + } + return true; + } + if ((type_aggregate == 4 /* TypeDesc::VEC4 */) || + (type_aggregate == 1 && type_arraylen == 4)) { + for (int i = 0; i < (derivatives ? 3 : 1); ++i) { + static_cast<ccl_private float *>(val)[i * 4 + 0] = fval[i].x; + static_cast<ccl_private float *>(val)[i * 4 + 1] = fval[i].y; + static_cast<ccl_private float *>(val)[i * 4 + 2] = fval[i].z; + static_cast<ccl_private float *>(val)[i * 4 + 3] = fval[i].w; + } + return true; + } + if ((type_aggregate == 1 /* TypeDesc::SCALAR */)) { + for (int i = 0; i < (derivatives ? 3 : 1); ++i) { + static_cast<ccl_private float *>(val)[i] = average(float4_to_float3(fval[i])); + } + return true; + } + } + + return false; +} +ccl_device_inline bool set_attribute_matrix(ccl_private const Transform &tfm, + TypeDesc type, + ccl_private void *val) +{ + const unsigned char type_basetype = type & 0xF; + const unsigned char type_aggregate = (type >> 8) & 0xF; + + if (type_basetype == 11 /* TypeDesc::FLOAT */ && type_aggregate == 16 /* TypeDesc::MATRIX44 */) { + copy_matrix(static_cast<ccl_private float *>(val), tfm); + return true; + } + + return false; +} + +ccl_device_inline bool get_background_attribute(KernelGlobals kg, + ccl_private ShaderData *sd, + DeviceString name, + TypeDesc type, + bool derivatives, + ccl_private void *val) +{ + if (name == DeviceStrings::u_path_ray_length) { + /* Ray Length */ + float f = sd->ray_length; + return set_attribute_float(f, type, derivatives, val); + } + + return false; +} + +ccl_device_inline bool get_object_attribute(KernelGlobals kg, + ccl_private ShaderData *sd, + const AttributeDescriptor &desc, + TypeDesc type, + bool derivatives, + ccl_private void *val) +{ + if (desc.type == NODE_ATTR_FLOAT) { + float fval[3]; +#ifdef __VOLUME__ + if (primitive_is_volume_attribute(sd, desc)) + fval[0] = primitive_volume_attribute_float(kg, sd, desc); + else +#endif + fval[0] = primitive_surface_attribute_float( + kg, sd, desc, derivatives ? &fval[1] : nullptr, derivatives ? &fval[2] : nullptr); + return set_attribute_float(fval, type, derivatives, val); + } + else if (desc.type == NODE_ATTR_FLOAT2) { + float2 fval[3]; +#ifdef __VOLUME__ + if (primitive_is_volume_attribute(sd, desc)) + return false; + else +#endif + fval[0] = primitive_surface_attribute_float2( + kg, sd, desc, derivatives ? &fval[1] : nullptr, derivatives ? &fval[2] : nullptr); + return set_attribute_float2(fval, type, derivatives, val); + } + else if (desc.type == NODE_ATTR_FLOAT3) { + float3 fval[3]; +#ifdef __VOLUME__ + if (primitive_is_volume_attribute(sd, desc)) + fval[0] = primitive_volume_attribute_float3(kg, sd, desc); + else +#endif + fval[0] = primitive_surface_attribute_float3( + kg, sd, desc, derivatives ? &fval[1] : nullptr, derivatives ? &fval[2] : nullptr); + return set_attribute_float3(fval, type, derivatives, val); + } + else if (desc.type == NODE_ATTR_FLOAT4 || desc.type == NODE_ATTR_RGBA) { + float4 fval[3]; +#ifdef __VOLUME__ + if (primitive_is_volume_attribute(sd, desc)) + fval[0] = primitive_volume_attribute_float4(kg, sd, desc); + else +#endif + fval[0] = primitive_surface_attribute_float4( + kg, sd, desc, derivatives ? &fval[1] : nullptr, derivatives ? &fval[2] : nullptr); + return set_attribute_float4(fval, type, derivatives, val); + } + else if (desc.type == NODE_ATTR_MATRIX) { + Transform tfm = primitive_attribute_matrix(kg, desc); + return set_attribute_matrix(tfm, type, val); + } + + return false; +} + +ccl_device_inline bool get_object_standard_attribute(KernelGlobals kg, + ccl_private ShaderData *sd, + DeviceString name, + TypeDesc type, + bool derivatives, + ccl_private void *val) +{ + /* Object attributes */ + if (name == DeviceStrings::u_object_location) { + float3 f = object_location(kg, sd); + return set_attribute_float3(f, type, derivatives, val); + } + else if (name == DeviceStrings::u_object_color) { + float3 f = object_color(kg, sd->object); + return set_attribute_float3(f, type, derivatives, val); + } + else if (name == DeviceStrings::u_object_alpha) { + float f = object_alpha(kg, sd->object); + return set_attribute_float(f, type, derivatives, val); + } + else if (name == DeviceStrings::u_object_index) { + float f = object_pass_id(kg, sd->object); + return set_attribute_float(f, type, derivatives, val); + } + else if (name == DeviceStrings::u_geom_dupli_generated) { + float3 f = object_dupli_generated(kg, sd->object); + return set_attribute_float3(f, type, derivatives, val); + } + else if (name == DeviceStrings::u_geom_dupli_uv) { + float3 f = object_dupli_uv(kg, sd->object); + return set_attribute_float3(f, type, derivatives, val); + } + else if (name == DeviceStrings::u_material_index) { + float f = shader_pass_id(kg, sd); + return set_attribute_float(f, type, derivatives, val); + } + else if (name == DeviceStrings::u_object_random) { + float f = object_random_number(kg, sd->object); + return set_attribute_float(f, type, derivatives, val); + } + + /* Particle attributes */ + else if (name == DeviceStrings::u_particle_index) { + int particle_id = object_particle_id(kg, sd->object); + float f = particle_index(kg, particle_id); + return set_attribute_float(f, type, derivatives, val); + } + else if (name == DeviceStrings::u_particle_random) { + int particle_id = object_particle_id(kg, sd->object); + float f = hash_uint2_to_float(particle_index(kg, particle_id), 0); + return set_attribute_float(f, type, derivatives, val); + } + + else if (name == DeviceStrings::u_particle_age) { + int particle_id = object_particle_id(kg, sd->object); + float f = particle_age(kg, particle_id); + return set_attribute_float(f, type, derivatives, val); + } + else if (name == DeviceStrings::u_particle_lifetime) { + int particle_id = object_particle_id(kg, sd->object); + float f = particle_lifetime(kg, particle_id); + return set_attribute_float(f, type, derivatives, val); + } + else if (name == DeviceStrings::u_particle_location) { + int particle_id = object_particle_id(kg, sd->object); + float3 f = particle_location(kg, particle_id); + return set_attribute_float3(f, type, derivatives, val); + } +#if 0 /* unsupported */ + else if (name == DeviceStrings::u_particle_rotation) { + int particle_id = object_particle_id(kg, sd->object); + float4 f = particle_rotation(kg, particle_id); + return set_attribute_float4(f, type, derivatives, val); + } +#endif + else if (name == DeviceStrings::u_particle_size) { + int particle_id = object_particle_id(kg, sd->object); + float f = particle_size(kg, particle_id); + return set_attribute_float(f, type, derivatives, val); + } + else if (name == DeviceStrings::u_particle_velocity) { + int particle_id = object_particle_id(kg, sd->object); + float3 f = particle_velocity(kg, particle_id); + return set_attribute_float3(f, type, derivatives, val); + } + else if (name == DeviceStrings::u_particle_angular_velocity) { + int particle_id = object_particle_id(kg, sd->object); + float3 f = particle_angular_velocity(kg, particle_id); + return set_attribute_float3(f, type, derivatives, val); + } + + /* Geometry attributes */ +#if 0 /* TODO */ + else if (name == DeviceStrings::u_geom_numpolyvertices) { + return false; + } + else if (name == DeviceStrings::u_geom_trianglevertices || + name == DeviceStrings::u_geom_polyvertices) { + return false; + } + else if (name == DeviceStrings::u_geom_name) { + return false; + } +#endif + else if (name == DeviceStrings::u_is_smooth) { + float f = ((sd->shader & SHADER_SMOOTH_NORMAL) != 0); + return set_attribute_float(f, type, derivatives, val); + } + +#ifdef __HAIR__ + /* Hair attributes */ + else if (name == DeviceStrings::u_is_curve) { + float f = (sd->type & PRIMITIVE_CURVE) != 0; + return set_attribute_float(f, type, derivatives, val); + } + else if (name == DeviceStrings::u_curve_thickness) { + float f = curve_thickness(kg, sd); + return set_attribute_float(f, type, derivatives, val); + } + else if (name == DeviceStrings::u_curve_tangent_normal) { + float3 f = curve_tangent_normal(kg, sd); + return set_attribute_float3(f, type, derivatives, val); + } + else if (name == DeviceStrings::u_curve_random) { + float f = curve_random(kg, sd); + return set_attribute_float(f, type, derivatives, val); + } +#endif + +#ifdef __POINTCLOUD__ + /* Point attributes */ + else if (name == DeviceStrings::u_is_point) { + float f = (sd->type & PRIMITIVE_POINT) != 0; + return set_attribute_float(f, type, derivatives, val); + } + else if (name == DeviceStrings::u_point_radius) { + float f = point_radius(kg, sd); + return set_attribute_float(f, type, derivatives, val); + } + else if (name == DeviceStrings::u_point_position) { + float3 f = point_position(kg, sd); + return set_attribute_float3(f, type, derivatives, val); + } + else if (name == DeviceStrings::u_point_random) { + float f = point_random(kg, sd); + return set_attribute_float(f, type, derivatives, val); + } +#endif + + else if (name == DeviceStrings::u_normal_map_normal) { + if (sd->type & PRIMITIVE_TRIANGLE) { + float3 f = triangle_smooth_normal_unnormalized(kg, sd, sd->Ng, sd->prim, sd->u, sd->v); + return set_attribute_float3(f, type, derivatives, val); + } + else { + return false; + } + } + + return get_background_attribute(kg, sd, name, type, derivatives, val); +} + +ccl_device_extern bool osl_get_attribute(ccl_private ShaderGlobals *sg, + int derivatives, + DeviceString object_name, + DeviceString name, + int array_lookup, + int index, + TypeDesc type, + ccl_private void *res) +{ + KernelGlobals kg = nullptr; + ccl_private ShaderData *const sd = static_cast<ccl_private ShaderData *>(sg->renderstate); + int object; + + if (object_name != DeviceStrings::_emptystring_) { + /* TODO: Get object index from name */ + return false; + } + else { + object = sd->object; + } + + const uint64_t id = name.hash(); + + const AttributeDescriptor desc = find_attribute(kg, object, sd->prim, sd->type, id); + if (desc.offset != ATTR_STD_NOT_FOUND) { + return get_object_attribute(kg, sd, desc, type, derivatives, res); + } + else { + return get_object_standard_attribute(kg, sd, name, type, derivatives, res); + } +} + +#if 0 +ccl_device_extern bool osl_bind_interpolated_param(ccl_private ShaderGlobals *sg, + DeviceString name, + long long type, + int userdata_has_derivs, + ccl_private void *userdata_data, + int symbol_has_derivs, + ccl_private void *symbol_data, + int symbol_data_size, + ccl_private void *userdata_initialized, + int userdata_index) +{ + return false; +} +#endif + +/* Noise */ + +#include "kernel/svm/noise.h" +#include "util/hash.h" + +ccl_device_extern uint osl_hash_ii(int x) +{ + return hash_uint(x); +} + +ccl_device_extern uint osl_hash_if(float x) +{ + return hash_uint(__float_as_uint(x)); +} + +ccl_device_extern uint osl_hash_iff(float x, float y) +{ + return hash_uint2(__float_as_uint(x), __float_as_uint(y)); +} + +ccl_device_extern uint osl_hash_iv(ccl_private const float3 *v) +{ + return hash_uint3(__float_as_uint(v->x), __float_as_uint(v->y), __float_as_uint(v->z)); +} + +ccl_device_extern uint osl_hash_ivf(ccl_private const float3 *v, float w) +{ + return hash_uint4( + __float_as_uint(v->x), __float_as_uint(v->y), __float_as_uint(v->z), __float_as_uint(w)); +} + +ccl_device_extern OSLNoiseOptions *osl_get_noise_options(ccl_private ShaderGlobals *sg) +{ + return nullptr; +} + +ccl_device_extern void osl_noiseparams_set_anisotropic(ccl_private OSLNoiseOptions *opt, + int anisotropic) +{ +} + +ccl_device_extern void osl_noiseparams_set_do_filter(ccl_private OSLNoiseOptions *opt, + int do_filter) +{ +} + +ccl_device_extern void osl_noiseparams_set_direction(ccl_private OSLNoiseOptions *opt, + float3 *direction) +{ +} + +ccl_device_extern void osl_noiseparams_set_bandwidth(ccl_private OSLNoiseOptions *opt, + float bandwidth) +{ +} + +ccl_device_extern void osl_noiseparams_set_impulses(ccl_private OSLNoiseOptions *opt, + float impulses) +{ +} + +#define OSL_NOISE_IMPL(name, op) \ + ccl_device_extern float name##_ff(float x) \ + { \ + return op##_1d(x); \ + } \ + ccl_device_extern float name##_fff(float x, float y) \ + { \ + return op##_2d(make_float2(x, y)); \ + } \ + ccl_device_extern float name##_fv(ccl_private const float3 *v) \ + { \ + return op##_3d(*v); \ + } \ + ccl_device_extern float name##_fvf(ccl_private const float3 *v, float w) \ + { \ + return op##_4d(make_float4(v->x, v->y, v->z, w)); \ + } \ + ccl_device_extern void name##_vf(ccl_private float3 *res, float x) \ + { \ + /* TODO: This is not correct. Really need to change the hash function inside the noise \ + * function to spit out a vector instead of a scalar. */ \ + const float n = name##_ff(x); \ + res->x = n; \ + res->y = n; \ + res->z = n; \ + } \ + ccl_device_extern void name##_vff(ccl_private float3 *res, float x, float y) \ + { \ + const float n = name##_fff(x, y); \ + res->x = n; \ + res->y = n; \ + res->z = n; \ + } \ + ccl_device_extern void name##_vv(ccl_private float3 *res, const float3 *v) \ + { \ + const float n = name##_fv(v); \ + res->x = n; \ + res->y = n; \ + res->z = n; \ + } \ + ccl_device_extern void name##_vvf(ccl_private float3 *res, const float3 *v, float w) \ + { \ + const float n = name##_fvf(v, w); \ + res->x = n; \ + res->y = n; \ + res->z = n; \ + } + +ccl_device_forceinline float hashnoise_1d(float p) +{ + const uint x = __float_as_uint(p); + return hash_uint(x) / static_cast<float>(~0u); +} +ccl_device_forceinline float hashnoise_2d(float2 p) +{ + const uint x = __float_as_uint(p.x); + const uint y = __float_as_uint(p.y); + return hash_uint2(x, y) / static_cast<float>(~0u); +} +ccl_device_forceinline float hashnoise_3d(float3 p) +{ + const uint x = __float_as_uint(p.x); + const uint y = __float_as_uint(p.y); + const uint z = __float_as_uint(p.z); + return hash_uint3(x, y, z) / static_cast<float>(~0u); +} +ccl_device_forceinline float hashnoise_4d(float4 p) +{ + const uint x = __float_as_uint(p.x); + const uint y = __float_as_uint(p.y); + const uint z = __float_as_uint(p.z); + const uint w = __float_as_uint(p.w); + return hash_uint4(x, y, z, w) / static_cast<float>(~0u); +} + +/* TODO: Implement all noise functions */ +OSL_NOISE_IMPL(osl_hashnoise, hashnoise) +OSL_NOISE_IMPL(osl_noise, noise) +OSL_NOISE_IMPL(osl_snoise, snoise) + +/* Texturing */ + +ccl_device_extern ccl_private OSLTextureOptions *osl_get_texture_options( + ccl_private ShaderGlobals *sg) +{ + return nullptr; +} + +ccl_device_extern void osl_texture_set_firstchannel(ccl_private OSLTextureOptions *opt, + int firstchannel) +{ +} + +ccl_device_extern void osl_texture_set_swrap_code(ccl_private OSLTextureOptions *opt, int mode) +{ +} + +ccl_device_extern void osl_texture_set_twrap_code(ccl_private OSLTextureOptions *opt, int mode) +{ +} + +ccl_device_extern void osl_texture_set_rwrap_code(ccl_private OSLTextureOptions *opt, int mode) +{ +} + +ccl_device_extern void osl_texture_set_stwrap_code(ccl_private OSLTextureOptions *opt, int mode) +{ +} + +ccl_device_extern void osl_texture_set_sblur(ccl_private OSLTextureOptions *opt, float blur) +{ +} + +ccl_device_extern void osl_texture_set_tblur(ccl_private OSLTextureOptions *opt, float blur) +{ +} + +ccl_device_extern void osl_texture_set_rblur(ccl_private OSLTextureOptions *opt, float blur) +{ +} + +ccl_device_extern void osl_texture_set_stblur(ccl_private OSLTextureOptions *opt, float blur) +{ +} + +ccl_device_extern void osl_texture_set_swidth(ccl_private OSLTextureOptions *opt, float width) +{ +} + +ccl_device_extern void osl_texture_set_twidth(ccl_private OSLTextureOptions *opt, float width) +{ +} + +ccl_device_extern void osl_texture_set_rwidth(ccl_private OSLTextureOptions *opt, float width) +{ +} + +ccl_device_extern void osl_texture_set_stwidth(ccl_private OSLTextureOptions *opt, float width) +{ +} + +ccl_device_extern void osl_texture_set_fill(ccl_private OSLTextureOptions *opt, float fill) +{ +} + +ccl_device_extern void osl_texture_set_time(ccl_private OSLTextureOptions *opt, float time) +{ +} + +ccl_device_extern void osl_texture_set_interp_code(ccl_private OSLTextureOptions *opt, int mode) +{ +} + +ccl_device_extern void osl_texture_set_subimage(ccl_private OSLTextureOptions *opt, int subimage) +{ +} + +ccl_device_extern void osl_texture_set_missingcolor_arena(ccl_private OSLTextureOptions *opt, + ccl_private float3 *color) +{ +} + +ccl_device_extern void osl_texture_set_missingcolor_alpha(ccl_private OSLTextureOptions *opt, + int nchannels, + float alpha) +{ +} + +ccl_device_extern bool osl_texture(ccl_private ShaderGlobals *sg, + DeviceString filename, + ccl_private void *texture_handle, + OSLTextureOptions *opt, + float s, + float t, + float dsdx, + float dtdx, + float dsdy, + float dtdy, + int nchannels, + ccl_private float *result, + ccl_private float *dresultdx, + ccl_private float *dresultdy, + ccl_private float *alpha, + ccl_private float *dalphadx, + ccl_private float *dalphady, + ccl_private void *errormessage) +{ + if (!texture_handle) { + return false; + } + + /* Only SVM textures are supported. */ + int id = static_cast<int>(reinterpret_cast<size_t>(texture_handle) - 1); + + const float4 rgba = kernel_tex_image_interp(nullptr, id, s, 1.0f - t); + + result[0] = rgba.x; + if (nchannels > 1) + result[1] = rgba.y; + if (nchannels > 2) + result[2] = rgba.z; + if (nchannels > 3) + result[3] = rgba.w; + + return true; +} + +ccl_device_extern bool osl_texture3d(ccl_private ShaderGlobals *sg, + DeviceString filename, + ccl_private void *texture_handle, + OSLTextureOptions *opt, + ccl_private const float3 *P, + ccl_private const float3 *dPdx, + ccl_private const float3 *dPdy, + ccl_private const float3 *dPdz, + int nchannels, + ccl_private float *result, + ccl_private float *dresultds, + ccl_private float *dresultdt, + ccl_private float *alpha, + ccl_private float *dalphadx, + ccl_private float *dalphady, + ccl_private void *errormessage) +{ + if (!texture_handle) { + return false; + } + + /* Only SVM textures are supported. */ + int id = static_cast<int>(reinterpret_cast<size_t>(texture_handle) - 1); + + const float4 rgba = kernel_tex_image_interp_3d(nullptr, id, *P, INTERPOLATION_NONE); + + result[0] = rgba.x; + if (nchannels > 1) + result[1] = rgba.y; + if (nchannels > 2) + result[2] = rgba.z; + if (nchannels > 3) + result[3] = rgba.w; + + return true; +} + +ccl_device_extern bool osl_environment(ccl_private ShaderGlobals *sg, + DeviceString filename, + ccl_private void *texture_handle, + OSLTextureOptions *opt, + ccl_private const float3 *R, + ccl_private const float3 *dRdx, + ccl_private const float3 *dRdy, + int nchannels, + ccl_private float *result, + ccl_private float *dresultds, + ccl_private float *dresultdt, + ccl_private float *alpha, + ccl_private float *dalphax, + ccl_private float *dalphay, + ccl_private void *errormessage) +{ + result[0] = 1.0f; + if (nchannels > 1) + result[1] = 0.0f; + if (nchannels > 2) + result[2] = 1.0f; + if (nchannels > 3) + result[3] = 1.0f; + + return false; +} + +ccl_device_extern bool osl_get_textureinfo(ccl_private ShaderGlobals *sg, + DeviceString filename, + ccl_private void *texture_handle, + DeviceString dataname, + int basetype, + int arraylen, + int aggegrate, + ccl_private void *data, + ccl_private void *errormessage) +{ + return false; +} + +ccl_device_extern bool osl_get_textureinfo_st(ccl_private ShaderGlobals *sg, + DeviceString filename, + ccl_private void *texture_handle, + float s, + float t, + DeviceString dataname, + int basetype, + int arraylen, + int aggegrate, + ccl_private void *data, + ccl_private void *errormessage) +{ + return osl_get_textureinfo( + sg, filename, texture_handle, dataname, basetype, arraylen, aggegrate, data, errormessage); +} + +/* Standard library */ + +#define OSL_OP_IMPL_II(name, op) \ + ccl_device_extern int name##_ii(int a) \ + { \ + return op(a); \ + } +#define OSL_OP_IMPL_IF(name, op) \ + ccl_device_extern int name##_if(float a) \ + { \ + return op(a); \ + } +#define OSL_OP_IMPL_FF(name, op) \ + ccl_device_extern float name##_ff(float a) \ + { \ + return op(a); \ + } +#define OSL_OP_IMPL_DFDF(name, op) \ + ccl_device_extern void name##_dfdf(ccl_private float *res, ccl_private const float *a) \ + { \ + for (int i = 0; i < 3; ++i) { \ + res[i] = op(a[i]); \ + } \ + } +#define OSL_OP_IMPL_DFDV(name, op) \ + ccl_device_extern void name##_dfdv(ccl_private float *res, ccl_private const float3 *a) \ + { \ + for (int i = 0; i < 3; ++i) { \ + res[i] = op(a[i]); \ + } \ + } +#define OSL_OP_IMPL_FV(name, op) \ + ccl_device_extern float name##_fv(ccl_private const float3 *a) \ + { \ + return op(*a); \ + } +#define OSL_OP_IMPL_VV(name, op) \ + ccl_device_extern void name##_vv(ccl_private float3 *res, ccl_private const float3 *a) \ + { \ + *res = op(*a); \ + } +#define OSL_OP_IMPL_VV_(name, op) \ + ccl_device_extern void name##_vv(ccl_private float3 *res, ccl_private const float3 *a) \ + { \ + res->x = op(a->x); \ + res->y = op(a->y); \ + res->z = op(a->z); \ + } +#define OSL_OP_IMPL_DVDV(name, op) \ + ccl_device_extern void name##_dvdv(ccl_private float3 *res, ccl_private const float3 *a) \ + { \ + for (int i = 0; i < 3; ++i) { \ + res[i] = op(a[i]); \ + } \ + } +#define OSL_OP_IMPL_DVDV_(name, op) \ + ccl_device_extern void name##_dvdv(ccl_private float3 *res, ccl_private const float3 *a) \ + { \ + for (int i = 0; i < 3; ++i) { \ + res[i].x = op(a[i].x); \ + res[i].y = op(a[i].y); \ + res[i].z = op(a[i].z); \ + } \ + } + +#define OSL_OP_IMPL_III(name, op) \ + ccl_device_extern int name##_iii(int a, int b) \ + { \ + return op(a, b); \ + } +#define OSL_OP_IMPL_FFF(name, op) \ + ccl_device_extern float name##_fff(float a, float b) \ + { \ + return op(a, b); \ + } +#define OSL_OP_IMPL_FVV(name, op) \ + ccl_device_extern float name##_fvv(ccl_private const float3 *a, ccl_private const float3 *b) \ + { \ + return op(*a, *b); \ + } +#define OSL_OP_IMPL_DFFDF(name, op) \ + ccl_device_extern void name##_dffdf( \ + ccl_private float *res, float a, ccl_private const float *b) \ + { \ + for (int i = 0; i < 3; ++i) { \ + res[i] = op(a, b[i]); \ + } \ + } +#define OSL_OP_IMPL_DFDFF(name, op) \ + ccl_device_extern void name##_dfdff( \ + ccl_private float *res, ccl_private const float *a, float b) \ + { \ + for (int i = 0; i < 3; ++i) { \ + res[i] = op(a[i], b); \ + } \ + } +#define OSL_OP_IMPL_DFDFDF(name, op) \ + ccl_device_extern void name##_dfdfdf( \ + ccl_private float *res, ccl_private const float *a, ccl_private const float *b) \ + { \ + for (int i = 0; i < 3; ++i) { \ + res[i] = op(a[i], b[i]); \ + } \ + } +#define OSL_OP_IMPL_DFVDV(name, op) \ + ccl_device_extern void name##_dfvdv( \ + ccl_private float *res, ccl_private const float3 *a, ccl_private const float3 *b) \ + { \ + for (int i = 0; i < 3; ++i) { \ + res[i] = op(a[0], b[i]); \ + } \ + } +#define OSL_OP_IMPL_DFDVV(name, op) \ + ccl_device_extern void name##_dfdvv( \ + ccl_private float *res, ccl_private const float3 *a, ccl_private const float3 *b) \ + { \ + for (int i = 0; i < 3; ++i) { \ + res[i] = op(a[i], b[0]); \ + } \ + } +#define OSL_OP_IMPL_DFDVDV(name, op) \ + ccl_device_extern void name##_dfdvdv( \ + ccl_private float *res, ccl_private const float3 *a, ccl_private const float3 *b) \ + { \ + for (int i = 0; i < 3; ++i) { \ + res[i] = op(a[i], b[i]); \ + } \ + } +#define OSL_OP_IMPL_VVF_(name, op) \ + ccl_device_extern void name##_vvf( \ + ccl_private float3 *res, ccl_private const float3 *a, float b) \ + { \ + res->x = op(a->x, b); \ + res->y = op(a->y, b); \ + res->z = op(a->z, b); \ + } +#define OSL_OP_IMPL_VVV(name, op) \ + ccl_device_extern void name##_vvv( \ + ccl_private float3 *res, ccl_private const float3 *a, ccl_private const float3 *b) \ + { \ + *res = op(*a, *b); \ + } +#define OSL_OP_IMPL_VVV_(name, op) \ + ccl_device_extern void name##_vvv( \ + ccl_private float3 *res, ccl_private const float3 *a, ccl_private const float3 *b) \ + { \ + res->x = op(a->x, b->x); \ + res->y = op(a->y, b->y); \ + res->z = op(a->z, b->z); \ + } +#define OSL_OP_IMPL_DVVDF_(name, op) \ + ccl_device_extern void name##_dvvdf( \ + ccl_private float3 *res, ccl_private const float3 *a, ccl_private const float *b) \ + { \ + for (int i = 0; i < 3; ++i) { \ + res[i].x = op(a[0].x, b[i]); \ + res[i].y = op(a[0].y, b[i]); \ + res[i].z = op(a[0].z, b[i]); \ + } \ + } +#define OSL_OP_IMPL_DVDVF_(name, op) \ + ccl_device_extern void name##_dvdvf( \ + ccl_private float3 *res, ccl_private const float3 *a, float b) \ + { \ + for (int i = 0; i < 3; ++i) { \ + res[i].x = op(a[i].x, b); \ + res[i].y = op(a[i].y, b); \ + res[i].z = op(a[i].z, b); \ + } \ + } +#define OSL_OP_IMPL_DVVDV(name, op) \ + ccl_device_extern void name##_dvvdv( \ + ccl_private float3 *res, ccl_private const float3 *a, ccl_private const float3 *b) \ + { \ + for (int i = 0; i < 3; ++i) { \ + res[i] = op(a[0], b[i]); \ + } \ + } +#define OSL_OP_IMPL_DVVDV_(name, op) \ + ccl_device_extern void name##_dvvdv( \ + ccl_private float3 *res, ccl_private const float3 *a, ccl_private const float3 *b) \ + { \ + for (int i = 0; i < 3; ++i) { \ + res[i].x = op(a[0].x, b[i].x); \ + res[i].y = op(a[0].y, b[i].y); \ + res[i].z = op(a[0].z, b[i].z); \ + } \ + } +#define OSL_OP_IMPL_DVDVV(name, op) \ + ccl_device_extern void name##_dvdvv( \ + ccl_private float3 *res, ccl_private const float3 *a, ccl_private const float3 *b) \ + { \ + for (int i = 0; i < 3; ++i) { \ + res[i] = op(a[i], b[0]); \ + } \ + } +#define OSL_OP_IMPL_DVDVV_(name, op) \ + ccl_device_extern void name##_dvdvv( \ + ccl_private float3 *res, ccl_private const float3 *a, ccl_private const float3 *b) \ + { \ + for (int i = 0; i < 3; ++i) { \ + res[i].x = op(a[i].x, b[0].x); \ + res[i].y = op(a[i].y, b[0].y); \ + res[i].z = op(a[i].z, b[0].z); \ + } \ + } +#define OSL_OP_IMPL_DVDVDF_(name, op) \ + ccl_device_extern void name##_dvdvdf( \ + ccl_private float3 *res, ccl_private const float3 *a, ccl_private const float *b) \ + { \ + for (int i = 0; i < 3; ++i) { \ + res[i].x = op(a[i].x, b[i]); \ + res[i].y = op(a[i].y, b[i]); \ + res[i].z = op(a[i].z, b[i]); \ + } \ + } +#define OSL_OP_IMPL_DVDVDV(name, op) \ + ccl_device_extern void name##_dvdvdv( \ + ccl_private float3 *res, ccl_private const float3 *a, ccl_private const float3 *b) \ + { \ + for (int i = 0; i < 3; ++i) { \ + res[i] = op(a[i], b[i]); \ + } \ + } +#define OSL_OP_IMPL_DVDVDV_(name, op) \ + ccl_device_extern void name##_dvdvdv( \ + ccl_private float3 *res, ccl_private const float3 *a, ccl_private const float3 *b) \ + { \ + for (int i = 0; i < 3; ++i) { \ + res[i].x = op(a[i].x, b[i].x); \ + res[i].y = op(a[i].y, b[i].y); \ + res[i].z = op(a[i].z, b[i].z); \ + } \ + } + +#define OSL_OP_IMPL_FFFF(name, op) \ + ccl_device_extern float name##_ffff(float a, float b, float c) \ + { \ + return op(a, b, c); \ + } +#define OSL_OP_IMPL_DFFFDF(name, op) \ + ccl_device_extern void name##_dfffdf( \ + ccl_private float *res, float a, float b, ccl_private const float *c) \ + { \ + for (int i = 0; i < 3; ++i) { \ + res[i] = op(a, b, c[i]); \ + } \ + } +#define OSL_OP_IMPL_DFFDFF(name, op) \ + ccl_device_extern void name##_dffdff( \ + ccl_private float *res, float a, ccl_private const float *b, float c) \ + { \ + for (int i = 0; i < 3; ++i) { \ + res[i] = op(a, b[i], c); \ + } \ + } +#define OSL_OP_IMPL_DFFDFDF(name, op) \ + ccl_device_extern void name##_dffdfdf( \ + ccl_private float *res, float a, ccl_private const float *b, ccl_private const float *c) \ + { \ + for (int i = 0; i < 3; ++i) { \ + res[i] = op(a, b[i], c[i]); \ + } \ + } + +#define OSL_OP_IMPL_DFDFFF(name, op) \ + ccl_device_extern void name##_dfdfff( \ + ccl_private float *res, ccl_private const float *a, float b, float c) \ + { \ + for (int i = 0; i < 3; ++i) { \ + res[i] = op(a[i], b, c); \ + } \ + } +#define OSL_OP_IMPL_DFDFFDF(name, op) \ + ccl_device_extern void name##_dfdffdf( \ + ccl_private float *res, ccl_private const float *a, float b, ccl_private const float *c) \ + { \ + for (int i = 0; i < 3; ++i) { \ + res[i] = op(a[i], b, c[i]); \ + } \ + } +#define OSL_OP_IMPL_DFDFDFF(name, op) \ + ccl_device_extern void name##_dfdfdff( \ + ccl_private float *res, ccl_private const float *a, ccl_private const float *b, float c) \ + { \ + for (int i = 0; i < 3; ++i) { \ + res[i] = op(a[i], b[i], c); \ + } \ + } +#define OSL_OP_IMPL_DFDFDFDF(name, op) \ + ccl_device_extern void name##_dfdfdfdf(ccl_private float *res, \ + ccl_private const float *a, \ + ccl_private const float *b, \ + ccl_private const float *c) \ + { \ + for (int i = 0; i < 3; ++i) { \ + res[i] = op(a[i], b[i], c[i]); \ + } \ + } + +#define OSL_OP_IMPL_XX(name, op) \ + OSL_OP_IMPL_FF(name, op) \ + OSL_OP_IMPL_DFDF(name, op) \ + OSL_OP_IMPL_VV_(name, op) \ + OSL_OP_IMPL_DVDV_(name, op) + +#define OSL_OP_IMPL_XXX(name, op) \ + OSL_OP_IMPL_FFF(name, op) \ + OSL_OP_IMPL_DFFDF(name, op) \ + OSL_OP_IMPL_DFDFF(name, op) \ + OSL_OP_IMPL_DFDFDF(name, op) \ + OSL_OP_IMPL_VVV_(name, op) \ + OSL_OP_IMPL_DVVDV_(name, op) \ + OSL_OP_IMPL_DVDVV_(name, op) \ + OSL_OP_IMPL_DVDVDV_(name, op) + +OSL_OP_IMPL_XX(osl_acos, acosf) +OSL_OP_IMPL_XX(osl_asin, asinf) +OSL_OP_IMPL_XX(osl_atan, atanf) +OSL_OP_IMPL_XXX(osl_atan2, atan2f) +OSL_OP_IMPL_XX(osl_cos, cosf) +OSL_OP_IMPL_XX(osl_sin, sinf) +OSL_OP_IMPL_XX(osl_tan, tanf) +OSL_OP_IMPL_XX(osl_cosh, coshf) +OSL_OP_IMPL_XX(osl_sinh, sinhf) +OSL_OP_IMPL_XX(osl_tanh, tanhf) + +ccl_device_forceinline int safe_divide(int a, int b) +{ + return (b != 0) ? a / b : 0; +} +ccl_device_forceinline int safe_modulo(int a, int b) +{ + return (b != 0) ? a % b : 0; +} + +OSL_OP_IMPL_III(osl_safe_div, safe_divide) +OSL_OP_IMPL_FFF(osl_safe_div, safe_divide) +OSL_OP_IMPL_III(osl_safe_mod, safe_modulo) + +ccl_device_extern void osl_sincos_fff(float a, ccl_private float *b, ccl_private float *c) +{ + sincos(a, b, c); +} +ccl_device_extern void osl_sincos_dfdff(ccl_private const float *a, + ccl_private float *b, + ccl_private float *c) +{ + for (int i = 0; i < 3; ++i) + sincos(a[i], b + i, c); +} +ccl_device_extern void osl_sincos_dffdf(ccl_private const float *a, + ccl_private float *b, + ccl_private float *c) +{ + for (int i = 0; i < 3; ++i) + sincos(a[i], b, c + i); +} +ccl_device_extern void osl_sincos_dfdfdf(ccl_private const float *a, + ccl_private float *b, + ccl_private float *c) +{ + for (int i = 0; i < 3; ++i) + sincos(a[i], b + i, c + i); +} +ccl_device_extern void osl_sincos_vvv(ccl_private const float3 *a, + ccl_private float3 *b, + ccl_private float3 *c) +{ + sincos(a->x, &b->x, &c->x); + sincos(a->y, &b->y, &c->y); + sincos(a->z, &b->z, &c->z); +} +ccl_device_extern void osl_sincos_dvdvv(ccl_private const float3 *a, + ccl_private float3 *b, + ccl_private float3 *c) +{ + for (int i = 0; i < 3; ++i) { + sincos(a[i].x, &b[i].x, &c->x); + sincos(a[i].y, &b[i].y, &c->y); + sincos(a[i].z, &b[i].z, &c->z); + } +} +ccl_device_extern void osl_sincos_dvvdv(ccl_private const float3 *a, + ccl_private float3 *b, + ccl_private float3 *c) +{ + for (int i = 0; i < 3; ++i) { + sincos(a[i].x, &b->x, &c[i].x); + sincos(a[i].y, &b->y, &c[i].y); + sincos(a[i].z, &b->z, &c[i].z); + } +} +ccl_device_extern void osl_sincos_dvdvdv(ccl_private const float3 *a, + ccl_private float3 *b, + ccl_private float3 *c) +{ + for (int i = 0; i < 3; ++i) { + sincos(a[i].x, &b[i].x, &c[i].x); + sincos(a[i].y, &b[i].y, &c[i].y); + sincos(a[i].z, &b[i].z, &c[i].z); + } +} + +OSL_OP_IMPL_XX(osl_log, logf) +OSL_OP_IMPL_XX(osl_log2, log2f) +OSL_OP_IMPL_XX(osl_log10, log10f) +OSL_OP_IMPL_XX(osl_exp, expf) +OSL_OP_IMPL_XX(osl_exp2, exp2f) +OSL_OP_IMPL_XX(osl_expm1, expm1f) +OSL_OP_IMPL_XX(osl_erf, erff) +OSL_OP_IMPL_XX(osl_erfc, erfcf) + +OSL_OP_IMPL_XXX(osl_pow, safe_powf) +OSL_OP_IMPL_VVF_(osl_pow, safe_powf) +OSL_OP_IMPL_DVVDF_(osl_pow, safe_powf) +OSL_OP_IMPL_DVDVF_(osl_pow, safe_powf) +OSL_OP_IMPL_DVDVDF_(osl_pow, safe_powf) + +OSL_OP_IMPL_XX(osl_sqrt, sqrtf) +OSL_OP_IMPL_XX(osl_inversesqrt, 1.0f / sqrtf) +OSL_OP_IMPL_XX(osl_cbrt, cbrtf) + +OSL_OP_IMPL_FF(osl_logb, logbf) +OSL_OP_IMPL_VV_(osl_logb, logbf) + +OSL_OP_IMPL_FF(osl_floor, floorf) +OSL_OP_IMPL_VV_(osl_floor, floorf) +OSL_OP_IMPL_FF(osl_ceil, ceilf) +OSL_OP_IMPL_VV_(osl_ceil, ceilf) +OSL_OP_IMPL_FF(osl_round, roundf) +OSL_OP_IMPL_VV_(osl_round, roundf) +OSL_OP_IMPL_FF(osl_trunc, truncf) +OSL_OP_IMPL_VV_(osl_trunc, truncf) + +ccl_device_forceinline float step_impl(float edge, float x) +{ + return x < edge ? 0.0f : 1.0f; +} + +OSL_OP_IMPL_FF(osl_sign, compatible_signf) +OSL_OP_IMPL_VV_(osl_sign, compatible_signf) +OSL_OP_IMPL_FFF(osl_step, step_impl) +OSL_OP_IMPL_VVV_(osl_step, step_impl) + +OSL_OP_IMPL_IF(osl_isnan, isnan) +OSL_OP_IMPL_IF(osl_isinf, isinf) +OSL_OP_IMPL_IF(osl_isfinite, isfinite) + +OSL_OP_IMPL_II(osl_abs, abs) +OSL_OP_IMPL_XX(osl_abs, fabsf) +OSL_OP_IMPL_II(osl_fabs, abs) +OSL_OP_IMPL_XX(osl_fabs, fabsf) +OSL_OP_IMPL_XXX(osl_fmod, safe_modulo) + +OSL_OP_IMPL_FFFF(osl_smoothstep, smoothstep) +OSL_OP_IMPL_DFFFDF(osl_smoothstep, smoothstep) +OSL_OP_IMPL_DFFDFF(osl_smoothstep, smoothstep) +OSL_OP_IMPL_DFFDFDF(osl_smoothstep, smoothstep) +OSL_OP_IMPL_DFDFFF(osl_smoothstep, smoothstep) +OSL_OP_IMPL_DFDFFDF(osl_smoothstep, smoothstep) +OSL_OP_IMPL_DFDFDFF(osl_smoothstep, smoothstep) +OSL_OP_IMPL_DFDFDFDF(osl_smoothstep, smoothstep) + +OSL_OP_IMPL_FVV(osl_dot, dot) +OSL_OP_IMPL_DFDVV(osl_dot, dot) +OSL_OP_IMPL_DFVDV(osl_dot, dot) +OSL_OP_IMPL_DFDVDV(osl_dot, dot) +OSL_OP_IMPL_VVV(osl_cross, cross) +OSL_OP_IMPL_DVDVV(osl_cross, cross) +OSL_OP_IMPL_DVVDV(osl_cross, cross) +OSL_OP_IMPL_DVDVDV(osl_cross, cross) +OSL_OP_IMPL_FV(osl_length, len) +OSL_OP_IMPL_DFDV(osl_length, len) +OSL_OP_IMPL_FVV(osl_distance, distance) +OSL_OP_IMPL_DFDVV(osl_distance, distance) +OSL_OP_IMPL_DFVDV(osl_distance, distance) +OSL_OP_IMPL_DFDVDV(osl_distance, distance) +OSL_OP_IMPL_VV(osl_normalize, safe_normalize) +OSL_OP_IMPL_DVDV(osl_normalize, safe_normalize) + +ccl_device_extern void osl_calculatenormal(ccl_private float3 *res, + ccl_private ShaderGlobals *sg, + ccl_private const float3 *p) +{ + if (sg->flipHandedness) + *res = cross(p[2], p[1]); + else + *res = cross(p[1], p[2]); +} + +ccl_device_extern float osl_area(ccl_private const float3 *p) +{ + return len(cross(p[2], p[1])); +} + +ccl_device_extern float osl_filterwidth_fdf(ccl_private const float *x) +{ + return sqrtf(x[1] * x[1] + x[2] * x[2]); +} + +ccl_device_extern void osl_filterwidth_vdv(ccl_private float *res, ccl_private const float *x) +{ + for (int i = 0; i < 3; ++i) + res[i] = osl_filterwidth_fdf(x + i); +} + +ccl_device_extern bool osl_raytype_bit(ccl_private ShaderGlobals *sg, int bit) +{ + return (sg->raytype & bit) != 0; +} diff --git a/intern/cycles/kernel/osl/services_optix.cu b/intern/cycles/kernel/osl/services_optix.cu new file mode 100644 index 00000000000..2a43a89a956 --- /dev/null +++ b/intern/cycles/kernel/osl/services_optix.cu @@ -0,0 +1,17 @@ +/* SPDX-License-Identifier: Apache-2.0 + * Copyright 2011-2022 Blender Foundation */ + +#define WITH_OSL + +// clang-format off +#include "kernel/device/optix/compat.h" +#include "kernel/device/optix/globals.h" + +#include "kernel/device/gpu/image.h" /* Texture lookup uses normal CUDA intrinsics. */ + +#include "kernel/osl/services_gpu.h" +// clang-format on + +extern "C" __device__ void __direct_callable__dummy_services() +{ +} diff --git a/intern/cycles/kernel/osl/types.h b/intern/cycles/kernel/osl/types.h index 46e06114360..717306a3d07 100644 --- a/intern/cycles/kernel/osl/types.h +++ b/intern/cycles/kernel/osl/types.h @@ -5,9 +5,53 @@ CCL_NAMESPACE_BEGIN +struct DeviceString { +#if defined(__KERNEL_GPU__) + /* Strings are represented by their hashes in CUDA and OptiX. */ + size_t str_; + + ccl_device_inline_method uint64_t hash() const + { + return str_; + } +#elif defined(OPENIMAGEIO_USTRING_H) + ustring str_; + + ccl_device_inline_method uint64_t hash() const + { + return str_.hash(); + } +#else + const char *str_; +#endif + + ccl_device_inline_method bool operator==(DeviceString b) const + { + return str_ == b.str_; + } + ccl_device_inline_method bool operator!=(DeviceString b) const + { + return str_ != b.str_; + } +}; + +ccl_device_inline DeviceString make_string(const char *str, size_t hash) +{ +#if defined(__KERNEL_GPU__) + (void)str; + return {hash}; +#elif defined(OPENIMAGEIO_USTRING_H) + (void)hash; + return {ustring(str)}; +#else + (void)hash; + return {str}; +#endif +} + /* Closure */ -enum ClosureTypeOSL { +enum OSLClosureType { OSL_CLOSURE_MUL_ID = -1, OSL_CLOSURE_ADD_ID = -2, @@ -17,4 +61,60 @@ enum ClosureTypeOSL { #include "closures_template.h" }; +struct OSLClosure { + OSLClosureType id; +}; + +struct ccl_align(8) OSLClosureMul : public OSLClosure +{ + packed_float3 weight; + ccl_private const OSLClosure *closure; +}; + +struct ccl_align(8) OSLClosureAdd : public OSLClosure +{ + ccl_private const OSLClosure *closureA; + ccl_private const OSLClosure *closureB; +}; + +struct ccl_align(8) OSLClosureComponent : public OSLClosure +{ + packed_float3 weight; +}; + +/* Globals */ + +struct ShaderGlobals { + packed_float3 P, dPdx, dPdy; + packed_float3 dPdz; + packed_float3 I, dIdx, dIdy; + packed_float3 N; + packed_float3 Ng; + float u, dudx, dudy; + float v, dvdx, dvdy; + packed_float3 dPdu, dPdv; + float time; + float dtime; + packed_float3 dPdtime; + packed_float3 Ps, dPsdx, dPsdy; + ccl_private void *renderstate; + ccl_private void *tracedata; + ccl_private void *objdata; + void *context; + void *renderer; + ccl_private void *object2common; + ccl_private void *shader2common; + ccl_private OSLClosure *Ci; + float surfacearea; + int raytype; + int flipHandedness; + int backfacing; +}; + +struct OSLNoiseOptions { +}; + +struct OSLTextureOptions { +}; + CCL_NAMESPACE_END diff --git a/intern/cycles/kernel/types.h b/intern/cycles/kernel/types.h index 24c5a6a4540..a6f8914a9b8 100644 --- a/intern/cycles/kernel/types.h +++ b/intern/cycles/kernel/types.h @@ -75,10 +75,14 @@ CCL_NAMESPACE_BEGIN #define __VOLUME__ /* Device specific features */ -#ifndef __KERNEL_GPU__ -# ifdef WITH_OSL -# define __OSL__ +#ifdef WITH_OSL +# define __OSL__ +# ifdef __KERNEL_OPTIX__ +/* Kernels with OSL support are built separately in OptiX and don't need SVM. */ +# undef __SVM__ # endif +#endif +#ifndef __KERNEL_GPU__ # ifdef WITH_PATH_GUIDING # define __PATH_GUIDING__ # endif @@ -917,9 +921,13 @@ typedef struct ccl_align(16) ShaderData float ray_dP; #ifdef __OSL__ +# ifdef __KERNEL_GPU__ + ccl_private uint8_t *osl_closure_pool; +# else const struct KernelGlobalsCPU *osl_globals; const struct IntegratorStateCPU *osl_path_state; const struct IntegratorShadowStateCPU *osl_shadow_path_state; +# endif #endif /* LCG state for closures that require additional random numbers. */ @@ -1529,6 +1537,9 @@ enum KernelFeatureFlag : uint32_t { /* Path guiding. */ KERNEL_FEATURE_PATH_GUIDING = (1U << 26U), + + /* OSL. */ + KERNEL_FEATURE_OSL = (1U << 27U), }; /* Shader node feature mask, to specialize shader evaluation for kernels. */ diff --git a/intern/cycles/scene/osl.cpp b/intern/cycles/scene/osl.cpp index 93839facdbe..3ea406b6935 100644 --- a/intern/cycles/scene/osl.cpp +++ b/intern/cycles/scene/osl.cpp @@ -38,16 +38,17 @@ OSL::TextureSystem *OSLShaderManager::ts_shared = NULL; int OSLShaderManager::ts_shared_users = 0; thread_mutex OSLShaderManager::ts_shared_mutex; -OSL::ShadingSystem *OSLShaderManager::ss_shared = NULL; -OSLRenderServices *OSLShaderManager::services_shared = NULL; +OSL::ErrorHandler OSLShaderManager::errhandler; +map<int, OSL::ShadingSystem *> OSLShaderManager::ss_shared; int OSLShaderManager::ss_shared_users = 0; thread_mutex OSLShaderManager::ss_shared_mutex; thread_mutex OSLShaderManager::ss_mutex; + int OSLCompiler::texture_shared_unique_id = 0; /* Shader Manager */ -OSLShaderManager::OSLShaderManager() +OSLShaderManager::OSLShaderManager(Device *device) : device_(device) { texture_system_init(); shading_system_init(); @@ -107,11 +108,12 @@ void OSLShaderManager::device_update_specific(Device *device, device_free(device, dscene, scene); - /* set texture system */ - scene->image_manager->set_osl_texture_system((void *)ts); + /* set texture system (only on CPU devices, since GPU devices cannot use OIIO) */ + if (device->info.type == DEVICE_CPU) { + scene->image_manager->set_osl_texture_system((void *)ts_shared); + } /* create shaders */ - OSLGlobals *og = (OSLGlobals *)device->get_cpu_osl_memory(); Shader *background_shader = scene->background->get_shader(scene); foreach (Shader *shader, scene->shaders) { @@ -125,22 +127,34 @@ void OSLShaderManager::device_update_specific(Device *device, * compile shaders alternating */ thread_scoped_lock lock(ss_mutex); - OSLCompiler compiler(this, services, ss, scene); - compiler.background = (shader == background_shader); - compiler.compile(og, shader); + device->foreach_device( + [this, scene, shader, background = (shader == background_shader)](Device *sub_device) { + OSLGlobals *og = (OSLGlobals *)sub_device->get_cpu_osl_memory(); + OSL::ShadingSystem *ss = ss_shared[sub_device->info.type]; + + OSLCompiler compiler(this, ss, scene); + compiler.background = background; + compiler.compile(og, shader); + }); if (shader->get_use_mis() && shader->has_surface_emission) scene->light_manager->tag_update(scene, LightManager::SHADER_COMPILED); } /* setup shader engine */ - og->ss = ss; - og->ts = ts; - og->services = services; - int background_id = scene->shader_manager->get_shader_id(background_shader); - og->background_state = og->surface_state[background_id & SHADER_MASK]; - og->use = true; + + device->foreach_device([background_id](Device *sub_device) { + OSLGlobals *og = (OSLGlobals *)sub_device->get_cpu_osl_memory(); + OSL::ShadingSystem *ss = ss_shared[sub_device->info.type]; + + og->ss = ss; + og->ts = ts_shared; + og->services = static_cast<OSLRenderServices *>(ss->renderer()); + + og->background_state = og->surface_state[background_id & SHADER_MASK]; + og->use = true; + }); foreach (Shader *shader, scene->shaders) shader->clear_modified(); @@ -148,8 +162,12 @@ void OSLShaderManager::device_update_specific(Device *device, update_flags = UPDATE_NONE; /* add special builtin texture types */ - services->textures.insert(ustring("@ao"), new OSLTextureHandle(OSLTextureHandle::AO)); - services->textures.insert(ustring("@bevel"), new OSLTextureHandle(OSLTextureHandle::BEVEL)); + for (const auto &[device_type, ss] : ss_shared) { + OSLRenderServices *services = static_cast<OSLRenderServices *>(ss->renderer()); + + services->textures.insert(ustring("@ao"), new OSLTextureHandle(OSLTextureHandle::AO)); + services->textures.insert(ustring("@bevel"), new OSLTextureHandle(OSLTextureHandle::BEVEL)); + } device_update_common(device, dscene, scene, progress); @@ -166,26 +184,35 @@ void OSLShaderManager::device_update_specific(Device *device, * is being freed after the Session is freed. */ thread_scoped_lock lock(ss_shared_mutex); - ss->optimize_all_groups(); + for (const auto &[device_type, ss] : ss_shared) { + ss->optimize_all_groups(); + } + } + + /* load kernels */ + if (!device->load_osl_kernels()) { + progress.set_error(device->error_message()); } } void OSLShaderManager::device_free(Device *device, DeviceScene *dscene, Scene *scene) { - OSLGlobals *og = (OSLGlobals *)device->get_cpu_osl_memory(); - device_free_common(device, dscene, scene); /* clear shader engine */ - og->use = false; - og->ss = NULL; - og->ts = NULL; - - og->surface_state.clear(); - og->volume_state.clear(); - og->displacement_state.clear(); - og->bump_state.clear(); - og->background_state.reset(); + device->foreach_device([](Device *sub_device) { + OSLGlobals *og = (OSLGlobals *)sub_device->get_cpu_osl_memory(); + + og->use = false; + og->ss = NULL; + og->ts = NULL; + + og->surface_state.clear(); + og->volume_state.clear(); + og->displacement_state.clear(); + og->bump_state.clear(); + og->background_state.reset(); + }); } void OSLShaderManager::texture_system_init() @@ -193,7 +220,7 @@ void OSLShaderManager::texture_system_init() /* create texture system, shared between different renders to reduce memory usage */ thread_scoped_lock lock(ts_shared_mutex); - if (ts_shared_users == 0) { + if (ts_shared_users++ == 0) { ts_shared = TextureSystem::create(true); ts_shared->attribute("automip", 1); @@ -203,24 +230,18 @@ void OSLShaderManager::texture_system_init() /* effectively unlimited for now, until we support proper mipmap lookups */ ts_shared->attribute("max_memory_MB", 16384); } - - ts = ts_shared; - ts_shared_users++; } void OSLShaderManager::texture_system_free() { /* shared texture system decrease users and destroy if no longer used */ thread_scoped_lock lock(ts_shared_mutex); - ts_shared_users--; - if (ts_shared_users == 0) { + if (--ts_shared_users == 0) { ts_shared->invalidate_all(true); OSL::TextureSystem::destroy(ts_shared); ts_shared = NULL; } - - ts = NULL; } void OSLShaderManager::shading_system_init() @@ -228,101 +249,105 @@ void OSLShaderManager::shading_system_init() /* create shading system, shared between different renders to reduce memory usage */ thread_scoped_lock lock(ss_shared_mutex); - if (ss_shared_users == 0) { - /* Must use aligned new due to concurrent hash map. */ - services_shared = util_aligned_new<OSLRenderServices>(ts_shared); + device_->foreach_device([](Device *sub_device) { + const DeviceType device_type = sub_device->info.type; - string shader_path = path_get("shader"); + if (ss_shared_users++ == 0 || ss_shared.find(device_type) == ss_shared.end()) { + /* Must use aligned new due to concurrent hash map. */ + OSLRenderServices *services = util_aligned_new<OSLRenderServices>(ts_shared, device_type); + + string shader_path = path_get("shader"); # ifdef _WIN32 - /* Annoying thing, Cycles stores paths in UTF-8 codepage, so it can - * operate with file paths with any character. This requires to use wide - * char functions, but OSL uses old fashioned ANSI functions which means: - * - * - We have to convert our paths to ANSI before passing to OSL - * - OSL can't be used when there's a multi-byte character in the path - * to the shaders folder. - */ - shader_path = string_to_ansi(shader_path); + /* Annoying thing, Cycles stores paths in UTF-8 codepage, so it can + * operate with file paths with any character. This requires to use wide + * char functions, but OSL uses old fashioned ANSI functions which means: + * + * - We have to convert our paths to ANSI before passing to OSL + * - OSL can't be used when there's a multi-byte character in the path + * to the shaders folder. + */ + shader_path = string_to_ansi(shader_path); # endif - ss_shared = new OSL::ShadingSystem(services_shared, ts_shared, &errhandler); - ss_shared->attribute("lockgeom", 1); - ss_shared->attribute("commonspace", "world"); - ss_shared->attribute("searchpath:shader", shader_path); - ss_shared->attribute("greedyjit", 1); - - VLOG_INFO << "Using shader search path: " << shader_path; - - /* our own ray types */ - static const char *raytypes[] = { - "camera", /* PATH_RAY_CAMERA */ - "reflection", /* PATH_RAY_REFLECT */ - "refraction", /* PATH_RAY_TRANSMIT */ - "diffuse", /* PATH_RAY_DIFFUSE */ - "glossy", /* PATH_RAY_GLOSSY */ - "singular", /* PATH_RAY_SINGULAR */ - "transparent", /* PATH_RAY_TRANSPARENT */ - "volume_scatter", /* PATH_RAY_VOLUME_SCATTER */ - - "shadow", /* PATH_RAY_SHADOW_OPAQUE */ - "shadow", /* PATH_RAY_SHADOW_TRANSPARENT */ - - "__unused__", /* PATH_RAY_NODE_UNALIGNED */ - "__unused__", /* PATH_RAY_MIS_SKIP */ - - "diffuse_ancestor", /* PATH_RAY_DIFFUSE_ANCESTOR */ - - /* Remaining irrelevant bits up to 32. */ - "__unused__", - "__unused__", - "__unused__", - "__unused__", - "__unused__", - "__unused__", - "__unused__", - "__unused__", - "__unused__", - "__unused__", - "__unused__", - "__unused__", - "__unused__", - "__unused__", - "__unused__", - "__unused__", - "__unused__", - "__unused__", - "__unused__", - }; - - const int nraytypes = sizeof(raytypes) / sizeof(raytypes[0]); - ss_shared->attribute("raytypes", TypeDesc(TypeDesc::STRING, nraytypes), raytypes); - - OSLRenderServices::register_closures(ss_shared); - - loaded_shaders.clear(); - } + OSL::ShadingSystem *ss = new OSL::ShadingSystem(services, ts_shared, &errhandler); + ss->attribute("lockgeom", 1); + ss->attribute("commonspace", "world"); + ss->attribute("searchpath:shader", shader_path); + ss->attribute("greedyjit", 1); + + VLOG_INFO << "Using shader search path: " << shader_path; + + /* our own ray types */ + static const char *raytypes[] = { + "camera", /* PATH_RAY_CAMERA */ + "reflection", /* PATH_RAY_REFLECT */ + "refraction", /* PATH_RAY_TRANSMIT */ + "diffuse", /* PATH_RAY_DIFFUSE */ + "glossy", /* PATH_RAY_GLOSSY */ + "singular", /* PATH_RAY_SINGULAR */ + "transparent", /* PATH_RAY_TRANSPARENT */ + "volume_scatter", /* PATH_RAY_VOLUME_SCATTER */ + + "shadow", /* PATH_RAY_SHADOW_OPAQUE */ + "shadow", /* PATH_RAY_SHADOW_TRANSPARENT */ + + "__unused__", /* PATH_RAY_NODE_UNALIGNED */ + "__unused__", /* PATH_RAY_MIS_SKIP */ + + "diffuse_ancestor", /* PATH_RAY_DIFFUSE_ANCESTOR */ + + /* Remaining irrelevant bits up to 32. */ + "__unused__", + "__unused__", + "__unused__", + "__unused__", + "__unused__", + "__unused__", + "__unused__", + "__unused__", + "__unused__", + "__unused__", + "__unused__", + "__unused__", + "__unused__", + "__unused__", + "__unused__", + "__unused__", + "__unused__", + "__unused__", + "__unused__", + }; + + const int nraytypes = sizeof(raytypes) / sizeof(raytypes[0]); + ss->attribute("raytypes", TypeDesc(TypeDesc::STRING, nraytypes), raytypes); + + OSLRenderServices::register_closures(ss); + + ss_shared[device_type] = ss; + } + }); - ss = ss_shared; - services = services_shared; - ss_shared_users++; + loaded_shaders.clear(); } void OSLShaderManager::shading_system_free() { /* shared shading system decrease users and destroy if no longer used */ thread_scoped_lock lock(ss_shared_mutex); - ss_shared_users--; - if (ss_shared_users == 0) { - delete ss_shared; - ss_shared = NULL; + device_->foreach_device([](Device * /*sub_device*/) { + if (--ss_shared_users == 0) { + for (const auto &[device_type, ss] : ss_shared) { + OSLRenderServices *services = static_cast<OSLRenderServices *>(ss->renderer()); - util_aligned_delete(services_shared); - services_shared = NULL; - } + delete ss; + + util_aligned_delete(services); + } - ss = NULL; - services = NULL; + ss_shared.clear(); + } + }); } bool OSLShaderManager::osl_compile(const string &inputfile, const string &outputfile) @@ -447,7 +472,9 @@ const char *OSLShaderManager::shader_load_filepath(string filepath) const char *OSLShaderManager::shader_load_bytecode(const string &hash, const string &bytecode) { - ss->LoadMemoryCompiledShader(hash.c_str(), bytecode.c_str()); + for (const auto &[device_type, ss] : ss_shared) { + ss->LoadMemoryCompiledShader(hash.c_str(), bytecode.c_str()); + } OSLShaderInfo info; @@ -599,11 +626,11 @@ OSLNode *OSLShaderManager::osl_node(ShaderGraph *graph, /* Graph Compiler */ -OSLCompiler::OSLCompiler(OSLShaderManager *manager, - OSLRenderServices *services, - OSL::ShadingSystem *ss, - Scene *scene) - : scene(scene), manager(manager), services(services), ss(ss) +OSLCompiler::OSLCompiler(OSLShaderManager *manager, OSL::ShadingSystem *ss, Scene *scene) + : scene(scene), + manager(manager), + services(static_cast<OSLRenderServices *>(ss->renderer())), + ss(ss) { current_type = SHADER_TYPE_SURFACE; current_shader = NULL; @@ -1105,7 +1132,12 @@ OSL::ShaderGroupRef OSLCompiler::compile_type(Shader *shader, ShaderGraph *graph { current_type = type; - OSL::ShaderGroupRef group = ss->ShaderGroupBegin(shader->name.c_str()); + string name = shader->name.string(); + /* Replace invalid characters. */ + for (size_t i; (i = name.find_first_of(" .,:;+-*/#")) != string::npos;) + name.replace(i, 1, "_"); + + OSL::ShaderGroupRef group = ss->ShaderGroupBegin(name); ShaderNode *output = graph->output(); ShaderNodeSet dependencies; diff --git a/intern/cycles/scene/osl.h b/intern/cycles/scene/osl.h index 76c6bd96ce1..c0e82a9dc8d 100644 --- a/intern/cycles/scene/osl.h +++ b/intern/cycles/scene/osl.h @@ -54,7 +54,7 @@ struct OSLShaderInfo { class OSLShaderManager : public ShaderManager { public: - OSLShaderManager(); + OSLShaderManager(Device *device); ~OSLShaderManager(); static void free_memory(); @@ -92,25 +92,22 @@ class OSLShaderManager : public ShaderManager { const std::string &bytecode_hash = "", const std::string &bytecode = ""); - protected: + private: void texture_system_init(); void texture_system_free(); void shading_system_init(); void shading_system_free(); - OSL::ShadingSystem *ss; - OSL::TextureSystem *ts; - OSLRenderServices *services; - OSL::ErrorHandler errhandler; + Device *device_; map<string, OSLShaderInfo> loaded_shaders; static OSL::TextureSystem *ts_shared; static thread_mutex ts_shared_mutex; static int ts_shared_users; - static OSL::ShadingSystem *ss_shared; - static OSLRenderServices *services_shared; + static OSL::ErrorHandler errhandler; + static map<int, OSL::ShadingSystem *> ss_shared; static thread_mutex ss_shared_mutex; static thread_mutex ss_mutex; static int ss_shared_users; @@ -123,10 +120,7 @@ class OSLShaderManager : public ShaderManager { class OSLCompiler { public: #ifdef WITH_OSL - OSLCompiler(OSLShaderManager *manager, - OSLRenderServices *services, - OSL::ShadingSystem *shadingsys, - Scene *scene); + OSLCompiler(OSLShaderManager *manager, OSL::ShadingSystem *shadingsys, Scene *scene); #endif void compile(OSLGlobals *og, Shader *shader); diff --git a/intern/cycles/scene/scene.cpp b/intern/cycles/scene/scene.cpp index 3a05bede7a3..d5be86e1db9 100644 --- a/intern/cycles/scene/scene.cpp +++ b/intern/cycles/scene/scene.cpp @@ -99,11 +99,8 @@ Scene::Scene(const SceneParams ¶ms_, Device *device) { memset((void *)&dscene.data, 0, sizeof(dscene.data)); - /* OSL only works on the CPU */ - if (device->info.has_osl) - shader_manager = ShaderManager::create(params.shadingsystem); - else - shader_manager = ShaderManager::create(SHADINGSYSTEM_SVM); + shader_manager = ShaderManager::create( + device->info.has_osl ? params.shadingsystem : SHADINGSYSTEM_SVM, device); light_manager = new LightManager(); geometry_manager = new GeometryManager(); diff --git a/intern/cycles/scene/shader.cpp b/intern/cycles/scene/shader.cpp index 56670c6e4e3..f176c19ec95 100644 --- a/intern/cycles/scene/shader.cpp +++ b/intern/cycles/scene/shader.cpp @@ -395,15 +395,16 @@ ShaderManager::~ShaderManager() { } -ShaderManager *ShaderManager::create(int shadingsystem) +ShaderManager *ShaderManager::create(int shadingsystem, Device *device) { ShaderManager *manager; (void)shadingsystem; /* Ignored when built without OSL. */ + (void)device; #ifdef WITH_OSL if (shadingsystem == SHADINGSYSTEM_OSL) { - manager = new OSLShaderManager(); + manager = new OSLShaderManager(device); } else #endif @@ -722,6 +723,10 @@ uint ShaderManager::get_kernel_features(Scene *scene) } } + if (use_osl()) { + kernel_features |= KERNEL_FEATURE_OSL; + } + return kernel_features; } diff --git a/intern/cycles/scene/shader.h b/intern/cycles/scene/shader.h index 2670776aca4..69b22d2ad19 100644 --- a/intern/cycles/scene/shader.h +++ b/intern/cycles/scene/shader.h @@ -170,7 +170,7 @@ class ShaderManager { UPDATE_NONE = 0u, }; - static ShaderManager *create(int shadingsystem); + static ShaderManager *create(int shadingsystem, Device *device); virtual ~ShaderManager(); virtual void reset(Scene *scene) = 0; diff --git a/intern/cycles/scene/shader_nodes.h b/intern/cycles/scene/shader_nodes.h index cc3a71a0697..a3a931bb0b3 100644 --- a/intern/cycles/scene/shader_nodes.h +++ b/intern/cycles/scene/shader_nodes.h @@ -1542,6 +1542,10 @@ class OSLNode final : public ShaderNode { { return true; } + virtual int get_feature() + { + return ShaderNode::get_feature() | KERNEL_FEATURE_NODE_RAYTRACE; + } virtual bool equals(const ShaderNode & /*other*/) { diff --git a/intern/cycles/util/defines.h b/intern/cycles/util/defines.h index 1969529eff0..d5be14c8eba 100644 --- a/intern/cycles/util/defines.h +++ b/intern/cycles/util/defines.h @@ -23,6 +23,7 @@ /* Leave inlining decisions to compiler for these, the inline keyword here * is not about performance but including function definitions in headers. */ # define ccl_device static inline +# define ccl_device_extern extern "C" # define ccl_device_noinline static inline # define ccl_device_noinline_cpu ccl_device_noinline diff --git a/intern/cycles/util/transform.h b/intern/cycles/util/transform.h index d7f95b7f296..0c39901a63c 100644 --- a/intern/cycles/util/transform.h +++ b/intern/cycles/util/transform.h @@ -196,14 +196,7 @@ ccl_device_inline Transform make_transform_frame(float3 N) return make_transform(dx.x, dx.y, dx.z, 0.0f, dy.x, dy.y, dy.z, 0.0f, N.x, N.y, N.z, 0.0f); } -#ifndef __KERNEL_GPU__ - -ccl_device_inline Transform transform_zero() -{ - Transform zero = {zero_float4(), zero_float4(), zero_float4()}; - return zero; -} - +#if !defined(__KERNEL_METAL__) ccl_device_inline Transform operator*(const Transform a, const Transform b) { float4 c_x = make_float4(b.x.x, b.y.x, b.z.x, 0.0f); @@ -218,6 +211,15 @@ ccl_device_inline Transform operator*(const Transform a, const Transform b) return t; } +#endif + +#ifndef __KERNEL_GPU__ + +ccl_device_inline Transform transform_zero() +{ + Transform zero = {zero_float4(), zero_float4(), zero_float4()}; + return zero; +} ccl_device_inline void print_transform(const char *label, const Transform &t) { |