griefith/test
1
0
Fork 0
Code Issues Pull Requests Actions 1 Packages Projects Releases Wiki Activity

Fix #143714: Cycles OptiX fails to render linear and ribbon curves together

Browse Source 
This case was not accounted for previously, but is now possible when
the new curves object has curves with type poly.

Pull Request: https://projects.blender.org/blender/blender/pulls/144087
This commit is contained in:
Brecht Van Lommel
2025-08-11 19:36:26 +02:00
committed by Brecht Van Lommel
parent 91e081417c
commit dce6269d1f
8 changed files with 149 additions and 77 deletions
Download Patch File Download Diff File Expand all files Collapse all files

5
intern/cycles/blender/curves.cpp
View File

@@ -675,6 +675,8 @@ void BlenderSync::sync_particle_hair(
}
}
}
hair->curve_shape = scene->params.hair_shape;
}
template<typename TypeInCycles, typename GetValueAtIndex>
@@ -1035,6 +1037,9 @@ void BlenderSync::sync_hair(Hair *hair, BObjectInfo &b_ob_info, bool motion, con
if (!b_types.is_empty() && b_types[0] == CURVE_TYPE_POLY) {
hair->curve_shape = CURVE_THICK_LINEAR;
}
else {
hair->curve_shape = scene->params.hair_shape;
}
}
void BlenderSync::sync_hair(BObjectInfo &b_ob_info, Hair *hair)

2
intern/cycles/blender/sync.h
View File

@@ -246,6 +246,8 @@ class BlenderSync {
bool use_adaptive_subdivision = false;
bool use_developer_ui;
CurveShapeType curve_shape = CURVE_RIBBON;
float dicing_rate;
int max_subdivisions;

170
intern/cycles/device/optix/device_impl.cpp
View File

@@ -349,16 +349,11 @@ bool OptiXDevice::load_kernels(const uint kernel_features)
pipeline_options.pipelineLaunchParamsVariableName = "kernel_params"; /* See globals.h */
pipeline_options.usesPrimitiveTypeFlags = OPTIX_PRIMITIVE_TYPE_FLAGS_TRIANGLE;
if (kernel_features & KERNEL_FEATURE_HAIR) {
if (kernel_features & KERNEL_FEATURE_HAIR_THICK) {
pipeline_options.usesPrimitiveTypeFlags |= OPTIX_PRIMITIVE_TYPE_FLAGS_ROUND_LINEAR |
OPTIX_PRIMITIVE_TYPE_FLAGS_ROUND_CATMULLROM;
}
else {
pipeline_options.usesPrimitiveTypeFlags |= OPTIX_PRIMITIVE_TYPE_FLAGS_CUSTOM;
}
if (kernel_features & KERNEL_FEATURE_HAIR_THICK) {
pipeline_options.usesPrimitiveTypeFlags |= OPTIX_PRIMITIVE_TYPE_FLAGS_ROUND_LINEAR |
OPTIX_PRIMITIVE_TYPE_FLAGS_ROUND_CATMULLROM;
}
if (kernel_features & KERNEL_FEATURE_POINTCLOUD) {
if (kernel_features & (KERNEL_FEATURE_HAIR_RIBBON | KERNEL_FEATURE_POINTCLOUD)) {
pipeline_options.usesPrimitiveTypeFlags |= OPTIX_PRIMITIVE_TYPE_FLAGS_CUSTOM;
}
@@ -435,68 +430,84 @@ bool OptiXDevice::load_kernels(const uint kernel_features)
group_descs[PG_HITV].hitgroup.moduleAH = optix_module;
group_descs[PG_HITV].hitgroup.entryFunctionNameAH = "__anyhit__kernel_optix_volume_test";
if (kernel_features & KERNEL_FEATURE_HAIR) {
if (kernel_features & KERNEL_FEATURE_HAIR_THICK) {
/* Built-in thick curve intersection. */
OptixBuiltinISOptions builtin_options = {};
builtin_options.builtinISModuleType = OPTIX_PRIMITIVE_TYPE_ROUND_CATMULLROM;
builtin_options.buildFlags = OPTIX_BUILD_FLAG_PREFER_FAST_TRACE |
OPTIX_BUILD_FLAG_ALLOW_COMPACTION |
OPTIX_BUILD_FLAG_ALLOW_UPDATE;
builtin_options.curveEndcapFlags = OPTIX_CURVE_ENDCAP_DEFAULT; /* Disable end-caps. */
builtin_options.usesMotionBlur = false;
OptixProgramGroupDesc ignore_desc = {};
ignore_desc.kind = OPTIX_PROGRAM_GROUP_KIND_HITGROUP;
ignore_desc.hitgroup.moduleCH = optix_module;
ignore_desc.hitgroup.entryFunctionNameCH = "__closesthit__kernel_optix_ignore";
ignore_desc.hitgroup.moduleAH = optix_module;
ignore_desc.hitgroup.entryFunctionNameAH = "__anyhit__kernel_optix_ignore";
if (kernel_features & KERNEL_FEATURE_HAIR_THICK) {
/* Built-in thick curve intersection. */
OptixBuiltinISOptions builtin_options = {};
builtin_options.builtinISModuleType = OPTIX_PRIMITIVE_TYPE_ROUND_CATMULLROM;
builtin_options.buildFlags = OPTIX_BUILD_FLAG_PREFER_FAST_TRACE |
OPTIX_BUILD_FLAG_ALLOW_COMPACTION | OPTIX_BUILD_FLAG_ALLOW_UPDATE;
builtin_options.curveEndcapFlags = OPTIX_CURVE_ENDCAP_DEFAULT; /* Disable end-caps. */
builtin_options.usesMotionBlur = false;
optix_assert(optixBuiltinISModuleGet(
context, &module_options, &pipeline_options, &builtin_options, &builtin_modules[0]));
group_descs[PG_HITD].hitgroup.moduleIS = builtin_modules[0];
group_descs[PG_HITD].hitgroup.entryFunctionNameIS = nullptr;
group_descs[PG_HITS].hitgroup.moduleIS = builtin_modules[0];
group_descs[PG_HITS].hitgroup.entryFunctionNameIS = nullptr;
if (pipeline_options.usesMotionBlur) {
builtin_options.usesMotionBlur = true;
optix_assert(optixBuiltinISModuleGet(
context, &module_options, &pipeline_options, &builtin_options, &builtin_modules[0]));
context, &module_options, &pipeline_options, &builtin_options, &builtin_modules[1]));
group_descs[PG_HITD].hitgroup.moduleIS = builtin_modules[0];
group_descs[PG_HITD].hitgroup.entryFunctionNameIS = nullptr;
group_descs[PG_HITS].hitgroup.moduleIS = builtin_modules[0];
group_descs[PG_HITS].hitgroup.entryFunctionNameIS = nullptr;
group_descs[PG_HITD_MOTION] = group_descs[PG_HITD];
group_descs[PG_HITD_MOTION].hitgroup.moduleIS = builtin_modules[1];
group_descs[PG_HITS_MOTION] = group_descs[PG_HITS];
group_descs[PG_HITS_MOTION].hitgroup.moduleIS = builtin_modules[1];
}
if (pipeline_options.usesMotionBlur) {
builtin_options.usesMotionBlur = true;
builtin_options.builtinISModuleType = OPTIX_PRIMITIVE_TYPE_ROUND_LINEAR;
builtin_options.usesMotionBlur = false;
optix_assert(optixBuiltinISModuleGet(
context, &module_options, &pipeline_options, &builtin_options, &builtin_modules[1]));
optix_assert(optixBuiltinISModuleGet(
context, &module_options, &pipeline_options, &builtin_options, &builtin_modules[2]));
group_descs[PG_HITD_MOTION] = group_descs[PG_HITD];
group_descs[PG_HITD_MOTION].hitgroup.moduleIS = builtin_modules[1];
group_descs[PG_HITS_MOTION] = group_descs[PG_HITS];
group_descs[PG_HITS_MOTION].hitgroup.moduleIS = builtin_modules[1];
}
group_descs[PG_HITD_CURVE_LINEAR] = group_descs[PG_HITD];
group_descs[PG_HITD_CURVE_LINEAR].hitgroup.moduleIS = builtin_modules[2];
group_descs[PG_HITS_CURVE_LINEAR] = group_descs[PG_HITS];
group_descs[PG_HITS_CURVE_LINEAR].hitgroup.moduleIS = builtin_modules[2];
group_descs[PG_HITV_CURVE_LINEAR] = ignore_desc;
group_descs[PG_HITL_CURVE_LINEAR] = ignore_desc;
builtin_options.builtinISModuleType = OPTIX_PRIMITIVE_TYPE_ROUND_LINEAR;
builtin_options.usesMotionBlur = false;
if (pipeline_options.usesMotionBlur) {
builtin_options.usesMotionBlur = true;
optix_assert(optixBuiltinISModuleGet(
context, &module_options, &pipeline_options, &builtin_options, &builtin_modules[2]));
context, &module_options, &pipeline_options, &builtin_options, &builtin_modules[3]));
group_descs[PG_HITD_CURVE_LINEAR] = group_descs[PG_HITD];
group_descs[PG_HITD_CURVE_LINEAR].hitgroup.moduleIS = builtin_modules[2];
group_descs[PG_HITS_CURVE_LINEAR] = group_descs[PG_HITS];
group_descs[PG_HITS_CURVE_LINEAR].hitgroup.moduleIS = builtin_modules[2];
if (pipeline_options.usesMotionBlur) {
builtin_options.usesMotionBlur = true;
optix_assert(optixBuiltinISModuleGet(
context, &module_options, &pipeline_options, &builtin_options, &builtin_modules[3]));
group_descs[PG_HITD_CURVE_LINEAR_MOTION] = group_descs[PG_HITD_CURVE_LINEAR];
group_descs[PG_HITD_CURVE_LINEAR_MOTION].hitgroup.moduleIS = builtin_modules[3];
group_descs[PG_HITS_CURVE_LINEAR_MOTION] = group_descs[PG_HITS_CURVE_LINEAR];
group_descs[PG_HITS_CURVE_LINEAR_MOTION].hitgroup.moduleIS = builtin_modules[3];
}
}
else {
/* Custom ribbon intersection. */
group_descs[PG_HITD].hitgroup.moduleIS = optix_module;
group_descs[PG_HITS].hitgroup.moduleIS = optix_module;
group_descs[PG_HITD].hitgroup.entryFunctionNameIS = "__intersection__curve_ribbon";
group_descs[PG_HITS].hitgroup.entryFunctionNameIS = "__intersection__curve_ribbon";
group_descs[PG_HITD_CURVE_LINEAR_MOTION] = group_descs[PG_HITD_CURVE_LINEAR];
group_descs[PG_HITD_CURVE_LINEAR_MOTION].hitgroup.moduleIS = builtin_modules[3];
group_descs[PG_HITS_CURVE_LINEAR_MOTION] = group_descs[PG_HITS_CURVE_LINEAR];
group_descs[PG_HITS_CURVE_LINEAR_MOTION].hitgroup.moduleIS = builtin_modules[3];
group_descs[PG_HITV_CURVE_LINEAR_MOTION] = ignore_desc;
group_descs[PG_HITL_CURVE_LINEAR_MOTION] = ignore_desc;
}
}
if (kernel_features & KERNEL_FEATURE_HAIR_RIBBON) {
/* Custom ribbon intersection. */
group_descs[PG_HITD_CURVE_RIBBON] = group_descs[PG_HITD];
group_descs[PG_HITD_CURVE_RIBBON].kind = OPTIX_PROGRAM_GROUP_KIND_HITGROUP;
group_descs[PG_HITD_CURVE_RIBBON].hitgroup.moduleIS = optix_module;
group_descs[PG_HITD_CURVE_RIBBON].hitgroup.entryFunctionNameIS =
"__intersection__curve_ribbon";
group_descs[PG_HITS_CURVE_RIBBON] = group_descs[PG_HITS];
group_descs[PG_HITS_CURVE_RIBBON].kind = OPTIX_PROGRAM_GROUP_KIND_HITGROUP;
group_descs[PG_HITS_CURVE_RIBBON].hitgroup.moduleIS = optix_module;
group_descs[PG_HITS_CURVE_RIBBON].hitgroup.entryFunctionNameIS =
"__intersection__curve_ribbon";
group_descs[PG_HITV_CURVE_RIBBON] = ignore_desc;
group_descs[PG_HITL_CURVE_RIBBON] = ignore_desc;
}
if (kernel_features & KERNEL_FEATURE_POINTCLOUD) {
group_descs[PG_HITD_POINTCLOUD] = group_descs[PG_HITD];
@@ -507,6 +518,8 @@ bool OptiXDevice::load_kernels(const uint kernel_features)
group_descs[PG_HITS_POINTCLOUD].kind = OPTIX_PROGRAM_GROUP_KIND_HITGROUP;
group_descs[PG_HITS_POINTCLOUD].hitgroup.moduleIS = optix_module;
group_descs[PG_HITS_POINTCLOUD].hitgroup.entryFunctionNameIS = "__intersection__point";
group_descs[PG_HITV_POINTCLOUD] = ignore_desc;
group_descs[PG_HITL_POINTCLOUD] = ignore_desc;
}
/* Add hit group for local intersections. */
@@ -647,6 +660,10 @@ bool OptiXDevice::load_kernels(const uint kernel_features)
trace_css = std::max(trace_css,
stack_size[PG_HITS_CURVE_LINEAR_MOTION].cssIS +
stack_size[PG_HITS_CURVE_LINEAR_MOTION].cssAH);
trace_css = std::max(
trace_css, stack_size[PG_HITD_CURVE_RIBBON].cssIS + stack_size[PG_HITD_CURVE_RIBBON].cssAH);
trace_css = std::max(
trace_css, stack_size[PG_HITS_CURVE_RIBBON].cssIS + stack_size[PG_HITS_CURVE_RIBBON].cssAH);
trace_css = std::max(
trace_css, stack_size[PG_HITD_POINTCLOUD].cssIS + stack_size[PG_HITD_POINTCLOUD].cssAH);
trace_css = std::max(
@@ -678,18 +695,32 @@ bool OptiXDevice::load_kernels(const uint kernel_features)
if (pipeline_options.usesMotionBlur) {
pipeline_groups.push_back(groups[PG_HITD_MOTION]);
pipeline_groups.push_back(groups[PG_HITS_MOTION]);
pipeline_groups.push_back(groups[PG_HITV_MOTION]);
pipeline_groups.push_back(groups[PG_HITL_MOTION]);
}
if (kernel_features & KERNEL_FEATURE_HAIR_THICK) {
pipeline_groups.push_back(groups[PG_HITD_CURVE_LINEAR]);
pipeline_groups.push_back(groups[PG_HITS_CURVE_LINEAR]);
pipeline_groups.push_back(groups[PG_HITV_CURVE_LINEAR]);
pipeline_groups.push_back(groups[PG_HITL_CURVE_LINEAR]);
if (pipeline_options.usesMotionBlur) {
pipeline_groups.push_back(groups[PG_HITD_CURVE_LINEAR_MOTION]);
pipeline_groups.push_back(groups[PG_HITS_CURVE_LINEAR_MOTION]);
pipeline_groups.push_back(groups[PG_HITV_CURVE_LINEAR_MOTION]);
pipeline_groups.push_back(groups[PG_HITL_CURVE_LINEAR_MOTION]);
}
}
if (kernel_features & KERNEL_FEATURE_HAIR_RIBBON) {
pipeline_groups.push_back(groups[PG_HITD_CURVE_RIBBON]);
pipeline_groups.push_back(groups[PG_HITS_CURVE_RIBBON]);
pipeline_groups.push_back(groups[PG_HITV_CURVE_RIBBON]);
pipeline_groups.push_back(groups[PG_HITL_CURVE_RIBBON]);
}
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_HITV_POINTCLOUD]);
pipeline_groups.push_back(groups[PG_HITL_POINTCLOUD]);
}
optix_assert(optixPipelineCreate(context,
@@ -738,6 +769,10 @@ bool OptiXDevice::load_kernels(const uint kernel_features)
pipeline_groups.push_back(groups[PG_HITS_CURVE_LINEAR_MOTION]);
}
}
if (kernel_features & KERNEL_FEATURE_HAIR_RIBBON) {
pipeline_groups.push_back(groups[PG_HITD_CURVE_RIBBON]);
pipeline_groups.push_back(groups[PG_HITS_CURVE_RIBBON]);
}
if (kernel_features & KERNEL_FEATURE_POINTCLOUD) {
pipeline_groups.push_back(groups[PG_HITD_POINTCLOUD]);
pipeline_groups.push_back(groups[PG_HITS_POINTCLOUD]);
@@ -1671,17 +1706,22 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
instance.visibilityMask = 0xFF;
}
if (ob->get_geometry()->is_hair() &&
static_cast<const Hair *>(ob->get_geometry())->curve_shape != CURVE_RIBBON)
{
if (static_cast<const Hair *>(ob->get_geometry())->curve_shape == CURVE_THICK_LINEAR) {
if (ob->get_geometry()->is_hair()) {
const Hair *hair = static_cast<const Hair *>(ob->get_geometry());
if (hair->curve_shape == CURVE_RIBBON) {
instance.sbtOffset = PG_HITD_CURVE_RIBBON - PG_HITD;
/* Also skip curve ribbons in local trace calls. */
instance.visibilityMask |= 4;
}
else if (hair->curve_shape == CURVE_THICK_LINEAR) {
instance.sbtOffset = PG_HITD_CURVE_LINEAR - PG_HITD;
if (pipeline_options.usesMotionBlur && ob->get_geometry()->has_motion_blur()) {
if (pipeline_options.usesMotionBlur && hair->has_motion_blur()) {
instance.sbtOffset = PG_HITD_CURVE_LINEAR_MOTION - PG_HITD;
}
}
else {
if (pipeline_options.usesMotionBlur && ob->get_geometry()->has_motion_blur()) {
if (pipeline_options.usesMotionBlur && hair->has_motion_blur()) {
/* Select between motion blur and non-motion blur built-in intersection module. */
instance.sbtOffset = PG_HITD_MOTION - PG_HITD;
}

22
intern/cycles/device/optix/device_impl.h
View File

@@ -20,6 +20,7 @@ struct KernelParamsOptiX;
/* List of OptiX program groups. */
enum {
/* Ray generation */
PG_RGEN_INTERSECT_CLOSEST,
PG_RGEN_INTERSECT_SHADOW,
PG_RGEN_INTERSECT_SUBSURFACE,
@@ -37,28 +38,47 @@ enum {
PG_RGEN_EVAL_BACKGROUND,
PG_RGEN_EVAL_CURVE_SHADOW_TRANSPARENCY,
PG_RGEN_INIT_FROM_CAMERA,
/* Miss */
PG_MISS,
/* Hit */
PG_HITD, /* Default hit group. */
PG_HITS, /* __SHADOW_RECORD_ALL__ hit group. */
PG_HITL, /* __BVH_LOCAL__ hit group (only used for triangles). */
PG_HITV, /* __VOLUME__ hit group. */
PG_HITD_MOTION,
PG_HITS_MOTION,
PG_HITL_MOTION,
PG_HITV_MOTION,
PG_HITD_CURVE_LINEAR,
PG_HITS_CURVE_LINEAR,
PG_HITV_CURVE_LINEAR,
PG_HITL_CURVE_LINEAR,
PG_HITD_CURVE_LINEAR_MOTION,
PG_HITS_CURVE_LINEAR_MOTION,
PG_HITV_CURVE_LINEAR_MOTION,
PG_HITL_CURVE_LINEAR_MOTION,
PG_HITD_CURVE_RIBBON,
PG_HITS_CURVE_RIBBON,
PG_HITV_CURVE_RIBBON,
PG_HITL_CURVE_RIBBON,
PG_HITD_POINTCLOUD,
PG_HITS_POINTCLOUD,
PG_HITV_POINTCLOUD,
PG_HITL_POINTCLOUD,
/* Callable */
PG_CALL_SVM_AO,
PG_CALL_SVM_BEVEL,
NUM_PROGRAM_GROUPS
};
static const int MISS_PROGRAM_GROUP_OFFSET = PG_MISS;
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 NUM_HIT_PROGRAM_GROUPS = 24;
static const int CALLABLE_PROGRAM_GROUPS_BASE = PG_CALL_SVM_AO;
static const int NUM_CALLABLE_PROGRAM_GROUPS = 2;

7
intern/cycles/kernel/device/optix/bvh.h
View File

@@ -50,6 +50,13 @@ extern "C" __global__ void __miss__kernel_optix_miss()
optixSetPayload_5(PRIMITIVE_NONE);
}
extern "C" __global__ void __anyhit__kernel_optix_ignore()
{
return optixIgnoreIntersection();
}
extern "C" __global__ void __closesthit__kernel_optix_ignore() {}
extern "C" __global__ void __anyhit__kernel_optix_local_hit()
{
#if defined(__HAIR__) || defined(__POINTCLOUD__)

3
intern/cycles/kernel/types.h
View File

@@ -87,8 +87,9 @@ CCL_NAMESPACE_BEGIN
/* BVH/sampling kernel features. */
#define KERNEL_FEATURE_POINTCLOUD (1U << 12U)
#define KERNEL_FEATURE_HAIR (1U << 13U)
#define KERNEL_FEATURE_HAIR_RIBBON (1U << 13U)
#define KERNEL_FEATURE_HAIR_THICK (1U << 14U)
#define KERNEL_FEATURE_HAIR (KERNEL_FEATURE_HAIR_RIBBON | KERNEL_FEATURE_HAIR_THICK)
#define KERNEL_FEATURE_OBJECT_MOTION (1U << 15U)
/* Denotes whether baking functionality is needed. */

4
intern/cycles/scene/geometry.cpp
View File

@@ -430,11 +430,7 @@ void GeometryManager::device_update_preprocess(Device *device, Scene *scene, Pro
}
if (geom->is_hair()) {
/* Set curve shape, still a global scene setting for now. */
Hair *hair = static_cast<Hair *>(geom);
if (hair->curve_shape != CURVE_THICK_LINEAR) {
hair->curve_shape = scene->params.hair_shape;
}
if (hair->need_update_rebuild) {
device_update_flags |= DEVICE_CURVE_DATA_NEEDS_REALLOC;

13
intern/cycles/scene/scene.cpp
View File

@@ -5,7 +5,9 @@
#include <cstdlib>
#include "bvh/bvh.h"
#include "device/device.h"
#include "scene/alembic.h"
#include "scene/background.h"
#include "scene/bake.h"
@@ -27,6 +29,7 @@
#include "scene/svm.h"
#include "scene/tables.h"
#include "scene/volume.h"
#include "session/session.h"
#include "util/guarded_allocator.h"
@@ -500,9 +503,6 @@ void Scene::update_kernel_features()
const bool use_motion = need_motion() == Scene::MotionType::MOTION_BLUR;
kernel_features |= KERNEL_FEATURE_PATH_TRACING;
if (params.hair_shape == CURVE_THICK || params.hair_shape == CURVE_THICK_LINEAR) {
kernel_features |= KERNEL_FEATURE_HAIR_THICK;
}
/* Track the max prim count in case the backend needs to rebuild BVHs or
* kernels to support different limits. */
@@ -532,9 +532,10 @@ void Scene::update_kernel_features()
kernel_features |= KERNEL_FEATURE_SHADOW_CATCHER;
}
if (geom->is_hair()) {
kernel_features |= KERNEL_FEATURE_HAIR;
kernel_max_prim_count = max(kernel_max_prim_count,
static_cast<Hair *>(geom)->num_segments());
const Hair *hair = static_cast<const Hair *>(geom);
kernel_features |= (hair->curve_shape == CURVE_RIBBON) ? KERNEL_FEATURE_HAIR_RIBBON :
KERNEL_FEATURE_HAIR_THICK;
kernel_max_prim_count = max(kernel_max_prim_count, hair->num_segments());
}
else if (geom->is_pointcloud()) {
kernel_features |= KERNEL_FEATURE_POINTCLOUD;