diff --git a/intern/cycles/blender/curves.cpp b/intern/cycles/blender/curves.cpp index 63315e2eacb..23d735ccb1e 100644 --- a/intern/cycles/blender/curves.cpp +++ b/intern/cycles/blender/curves.cpp @@ -675,6 +675,8 @@ void BlenderSync::sync_particle_hair( } } } + + hair->curve_shape = scene->params.hair_shape; } template @@ -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) diff --git a/intern/cycles/blender/sync.h b/intern/cycles/blender/sync.h index fd716c58b5a..a9e18c8304b 100644 --- a/intern/cycles/blender/sync.h +++ b/intern/cycles/blender/sync.h @@ -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; diff --git a/intern/cycles/device/optix/device_impl.cpp b/intern/cycles/device/optix/device_impl.cpp index d468fc9d7f8..b3b3faa4ae6 100644 --- a/intern/cycles/device/optix/device_impl.cpp +++ b/intern/cycles/device/optix/device_impl.cpp @@ -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(ob->get_geometry())->curve_shape != CURVE_RIBBON) - { - if (static_cast(ob->get_geometry())->curve_shape == CURVE_THICK_LINEAR) { + if (ob->get_geometry()->is_hair()) { + const Hair *hair = static_cast(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; } diff --git a/intern/cycles/device/optix/device_impl.h b/intern/cycles/device/optix/device_impl.h index a1570310e28..197ca9196f1 100644 --- a/intern/cycles/device/optix/device_impl.h +++ b/intern/cycles/device/optix/device_impl.h @@ -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; diff --git a/intern/cycles/kernel/device/optix/bvh.h b/intern/cycles/kernel/device/optix/bvh.h index de253ca5dd3..40905abe3ee 100644 --- a/intern/cycles/kernel/device/optix/bvh.h +++ b/intern/cycles/kernel/device/optix/bvh.h @@ -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__) diff --git a/intern/cycles/kernel/types.h b/intern/cycles/kernel/types.h index bf4a2cefe2c..ea70caa921b 100644 --- a/intern/cycles/kernel/types.h +++ b/intern/cycles/kernel/types.h @@ -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. */ diff --git a/intern/cycles/scene/geometry.cpp b/intern/cycles/scene/geometry.cpp index 828bdac0778..3c878d493f4 100644 --- a/intern/cycles/scene/geometry.cpp +++ b/intern/cycles/scene/geometry.cpp @@ -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(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; diff --git a/intern/cycles/scene/scene.cpp b/intern/cycles/scene/scene.cpp index bd7ea76a8d4..b4c8cd6cd08 100644 --- a/intern/cycles/scene/scene.cpp +++ b/intern/cycles/scene/scene.cpp @@ -5,7 +5,9 @@ #include #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(geom)->num_segments()); + const Hair *hair = static_cast(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;