From 6487395fa560336e56817239e4e93bfd62e17086 Mon Sep 17 00:00:00 2001 From: Patrick Mours Date: Tue, 29 Jul 2025 17:05:01 +0200 Subject: [PATCH] Cycles: Add linear curve shape Add new "Linear 3D Curves" option in the Curves panel in the render properties. This renders curves as linear segments rather than smooth curves, for faster render time at the cost of accuracy. On NVIDIA Blackwell GPUs, this can give a 6x speedup compared to smooth curves, due to hardware acceleration. On NVIDIA Ada there is still a 3x speedup, and CPU and other GPU backends will also render this faster. A difference with smooth curves is that these have end caps, as this was simpler to implement and they are usually helpful anyway. In the future this functionality will also be used to properly support the CURVE_TYPE_POLY on the new curves object. Pull Request: https://projects.blender.org/blender/blender/pulls/139735 --- intern/cycles/blender/addon/properties.py | 11 +- intern/cycles/blender/curves.cpp | 8 + intern/cycles/bvh/embree.cpp | 50 ++-- intern/cycles/device/metal/bvh.mm | 48 ++-- intern/cycles/device/optix/device_impl.cpp | 159 ++++++++++--- intern/cycles/device/optix/device_impl.h | 6 +- intern/cycles/kernel/bvh/shadow_all.h | 4 +- intern/cycles/kernel/bvh/traversal.h | 4 +- .../closure/bsdf_principled_hair_chiang.h | 3 +- .../closure/bsdf_principled_hair_huang.h | 2 +- intern/cycles/kernel/device/cpu/bvh.h | 3 +- .../cycles/kernel/device/metal/kernel.metal | 27 ++- intern/cycles/kernel/device/oneapi/kernel.cpp | 2 +- intern/cycles/kernel/device/optix/bvh.h | 19 -- intern/cycles/kernel/geom/curve_intersect.h | 223 ++++++++++++++++-- intern/cycles/kernel/geom/shader_data.h | 8 +- intern/cycles/kernel/types.h | 3 + intern/cycles/scene/geometry.cpp | 4 +- intern/cycles/scene/hair.cpp | 9 +- intern/cycles/scene/scene.cpp | 2 +- .../cycles_renders/hair_linear_close_up.png | 3 + .../eevee_renders/hair_linear_close_up.png | 3 + .../render/hair/hair_linear_close_up.blend | 3 + .../hair_linear_close_up.png | 3 + .../hair_linear_close_up.png | 3 + .../hair_linear_close_up.png | 3 + 26 files changed, 480 insertions(+), 133 deletions(-) create mode 100644 tests/files/render/hair/cycles_renders/hair_linear_close_up.png create mode 100644 tests/files/render/hair/eevee_renders/hair_linear_close_up.png create mode 100644 tests/files/render/hair/hair_linear_close_up.blend create mode 100644 tests/files/render/hair/storm_hydra_renders/hair_linear_close_up.png create mode 100644 tests/files/render/hair/storm_usd_renders/hair_linear_close_up.png create mode 100644 tests/files/render/hair/workbench_renders/hair_linear_close_up.png diff --git a/intern/cycles/blender/addon/properties.py b/intern/cycles/blender/addon/properties.py index 1f275f08c26..2f8fb8472d5 100644 --- a/intern/cycles/blender/addon/properties.py +++ b/intern/cycles/blender/addon/properties.py @@ -61,8 +61,15 @@ enum_filter_types = ( ) enum_curve_shape = ( - ('RIBBONS', "Rounded Ribbons", "Render curves as flat ribbons with rounded normals, for fast rendering"), - ('THICK', "3D Curves", "Render curves as circular 3D geometry, for accurate results when viewing closely"), + ('RIBBONS', + "Rounded Ribbons", + "Render curves as flat ribbons with rounded normals, for fast rendering"), + ('THICK', + "3D Curves", + "Render curves as circular 3D geometry, for accurate results when viewing closely"), + ('THICK_LINEAR', + "Linear 3D Curves", + "Render curves as circular 3D geometry, with linear interpolation between control points, for fast rendering"), ) enum_use_layer_samples = ( diff --git a/intern/cycles/blender/curves.cpp b/intern/cycles/blender/curves.cpp index 5dc6696f54b..63315e2eacb 100644 --- a/intern/cycles/blender/curves.cpp +++ b/intern/cycles/blender/curves.cpp @@ -1029,6 +1029,12 @@ void BlenderSync::sync_hair(Hair *hair, BObjectInfo &b_ob_info, bool motion, con else { export_hair_curves(scene, hair, b_curves, need_motion, motion_scale); } + + const blender::VArray b_types = b_curves.curve_types(); + /* This does not handle cases where the curve type is not the same across all curves */ + if (!b_types.is_empty() && b_types[0] == CURVE_TYPE_POLY) { + hair->curve_shape = CURVE_THICK_LINEAR; + } } void BlenderSync::sync_hair(BObjectInfo &b_ob_info, Hair *hair) @@ -1068,6 +1074,8 @@ void BlenderSync::sync_hair(BObjectInfo &b_ob_info, Hair *hair) hair->attributes.update(std::move(new_hair.attributes)); + hair->curve_shape = new_hair.curve_shape; + /* tag update */ /* Compares curve_keys rather than strands in order to handle quick hair diff --git a/intern/cycles/bvh/embree.cpp b/intern/cycles/bvh/embree.cpp index 2b6d372f578..c4296905c2d 100644 --- a/intern/cycles/bvh/embree.cpp +++ b/intern/cycles/bvh/embree.cpp @@ -459,22 +459,34 @@ void pack_motion_verts(const size_t num_curves, const Hair *hair, const T *verts, const float *curve_radius, - float4 *rtc_verts) + float4 *rtc_verts, + CurveShapeType curve_shape) { for (size_t j = 0; j < num_curves; ++j) { const Hair::Curve c = hair->get_curve(j); int fk = c.first_key; - int k = 1; - for (; k < c.num_keys + 1; ++k, ++fk) { - rtc_verts[k].x = verts[fk].x; - rtc_verts[k].y = verts[fk].y; - rtc_verts[k].z = verts[fk].z; - rtc_verts[k].w = curve_radius[fk]; + + if (curve_shape == CURVE_THICK_LINEAR) { + for (int k = 0; k < c.num_keys; ++k, ++fk) { + rtc_verts[k].x = verts[fk].x; + rtc_verts[k].y = verts[fk].y; + rtc_verts[k].z = verts[fk].z; + rtc_verts[k].w = curve_radius[fk]; + } + rtc_verts += c.num_keys; + } + else { + for (int k = 1; k < c.num_keys + 1; ++k, ++fk) { + rtc_verts[k].x = verts[fk].x; + rtc_verts[k].y = verts[fk].y; + rtc_verts[k].z = verts[fk].z; + rtc_verts[k].w = curve_radius[fk]; + } + /* Duplicate Embree's Catmull-Rom spline CVs at the start and end of each curve. */ + rtc_verts[0] = rtc_verts[1]; + rtc_verts[c.num_keys + 1] = rtc_verts[c.num_keys]; + rtc_verts += c.num_keys + 2; } - /* Duplicate Embree's Catmull-Rom spline CVs at the start and end of each curve. */ - rtc_verts[0] = rtc_verts[1]; - rtc_verts[k] = rtc_verts[k - 1]; - rtc_verts += c.num_keys + 2; } } @@ -537,12 +549,14 @@ void BVHEmbree::set_curve_vertex_buffer(RTCGeometry geom_id, const Hair *hair, c const size_t num_curves = hair->num_curves(); if (t == t_mid || attr_mP == nullptr) { const float3 *verts = hair->get_curve_keys().data(); - pack_motion_verts(num_curves, hair, verts, curve_radius, rtc_verts); + pack_motion_verts( + num_curves, hair, verts, curve_radius, rtc_verts, hair->curve_shape); } else { const int t_ = (t > t_mid) ? (t - 1) : t; const float4 *verts = &attr_mP->data_float4()[t_ * num_keys]; - pack_motion_verts(num_curves, hair, verts, curve_radius, rtc_verts); + pack_motion_verts( + num_curves, hair, verts, curve_radius, rtc_verts, hair->curve_shape); } } @@ -681,7 +695,9 @@ void BVHEmbree::add_curves(const Object *ob, const Hair *hair, const int i) num_segments += c.num_segments(); } - const enum RTCGeometryType type = (hair->curve_shape == CURVE_RIBBON ? + const enum RTCGeometryType type = (hair->curve_shape == CURVE_THICK_LINEAR ? + RTC_GEOMETRY_TYPE_ROUND_LINEAR_CURVE : + hair->curve_shape == CURVE_RIBBON ? RTC_GEOMETRY_TYPE_FLAT_CATMULL_ROM_CURVE : RTC_GEOMETRY_TYPE_ROUND_CATMULL_ROM_CURVE); @@ -711,8 +727,10 @@ void BVHEmbree::add_curves(const Object *ob, const Hair *hair, const int i) const Hair::Curve c = hair->get_curve(j); for (size_t k = 0; k < c.num_segments(); ++k) { rtc_indices[rtc_index] = c.first_key + k; - /* Room for extra CVs at Catmull-Rom splines. */ - rtc_indices[rtc_index] += j * 2; + if (hair->curve_shape != CURVE_THICK_LINEAR) { + /* Room for extra CVs at Catmull-Rom splines. */ + rtc_indices[rtc_index] += j * 2; + } ++rtc_index; } diff --git a/intern/cycles/device/metal/bvh.mm b/intern/cycles/device/metal/bvh.mm index 19c17ac7e64..3506d9ab579 100644 --- a/intern/cycles/device/metal/bvh.mm +++ b/intern/cycles/device/metal/bvh.mm @@ -445,8 +445,10 @@ bool BVHMetal::build_BLAS_hair(Progress &progress, int segCount = curve.num_segments(); int firstKey = curve.first_key; uint64_t idxBase = cpData.size(); - cpData.push_back(keys[firstKey]); - radiusData.push_back(radiuses[firstKey]); + if (hair->curve_shape != CURVE_THICK_LINEAR) { + cpData.push_back(keys[firstKey]); + radiusData.push_back(radiuses[firstKey]); + } for (int s = 0; s < segCount; ++s) { if (step == 0) { idxData.push_back(idxBase + s); @@ -455,9 +457,11 @@ bool BVHMetal::build_BLAS_hair(Progress &progress, radiusData.push_back(radiuses[firstKey + s]); } cpData.push_back(keys[firstKey + curve.num_keys - 1]); - cpData.push_back(keys[firstKey + curve.num_keys - 1]); - radiusData.push_back(radiuses[firstKey + curve.num_keys - 1]); radiusData.push_back(radiuses[firstKey + curve.num_keys - 1]); + if (hair->curve_shape != CURVE_THICK_LINEAR) { + cpData.push_back(keys[firstKey + curve.num_keys - 1]); + radiusData.push_back(radiuses[firstKey + curve.num_keys - 1]); + } } } @@ -502,11 +506,17 @@ bool BVHMetal::build_BLAS_hair(Progress &progress, geomDescCrv.radiusStride = sizeof(float); geomDescCrv.radiusFormat = MTLAttributeFormatFloat; geomDescCrv.segmentCount = idxData.size(); - geomDescCrv.segmentControlPointCount = 4; + geomDescCrv.segmentControlPointCount = (hair->curve_shape == CURVE_THICK_LINEAR) ? 2 : 4; geomDescCrv.curveType = (hair->curve_shape == CURVE_RIBBON) ? MTLCurveTypeFlat : MTLCurveTypeRound; - geomDescCrv.curveBasis = MTLCurveBasisCatmullRom; - geomDescCrv.curveEndCaps = MTLCurveEndCapsDisk; + if (hair->curve_shape == CURVE_THICK_LINEAR) { + geomDescCrv.curveBasis = MTLCurveBasisLinear; + geomDescCrv.curveEndCaps = MTLCurveEndCapsSphere; + } + else { + geomDescCrv.curveBasis = MTLCurveBasisCatmullRom; + geomDescCrv.curveEndCaps = MTLCurveEndCapsDisk; + } geomDescCrv.indexType = MTLIndexTypeUInt32; geomDescCrv.indexBuffer = idxBuffer; geomDescCrv.intersectionFunctionTableOffset = 1; @@ -537,18 +547,22 @@ bool BVHMetal::build_BLAS_hair(Progress &progress, const Hair::Curve curve = hair->get_curve(c); int segCount = curve.num_segments(); int firstKey = curve.first_key; - radiusData.push_back(radiuses[firstKey]); uint64_t idxBase = cpData.size(); - cpData.push_back(keys[firstKey]); + if (hair->curve_shape != CURVE_THICK_LINEAR) { + cpData.push_back(keys[firstKey]); + radiusData.push_back(radiuses[firstKey]); + } for (int s = 0; s < segCount; ++s) { idxData.push_back(idxBase + s); cpData.push_back(keys[firstKey + s]); radiusData.push_back(radiuses[firstKey + s]); } cpData.push_back(keys[firstKey + curve.num_keys - 1]); - cpData.push_back(keys[firstKey + curve.num_keys - 1]); - radiusData.push_back(radiuses[firstKey + curve.num_keys - 1]); radiusData.push_back(radiuses[firstKey + curve.num_keys - 1]); + if (hair->curve_shape != CURVE_THICK_LINEAR) { + cpData.push_back(keys[firstKey + curve.num_keys - 1]); + radiusData.push_back(radiuses[firstKey + curve.num_keys - 1]); + } } /* Allocate and populate MTLBuffers for geometry. */ @@ -571,11 +585,17 @@ bool BVHMetal::build_BLAS_hair(Progress &progress, geomDescCrv.controlPointFormat = MTLAttributeFormatFloat3; geomDescCrv.controlPointBufferOffset = 0; geomDescCrv.segmentCount = idxData.size(); - geomDescCrv.segmentControlPointCount = 4; + geomDescCrv.segmentControlPointCount = (hair->curve_shape == CURVE_THICK_LINEAR) ? 2 : 4; geomDescCrv.curveType = (hair->curve_shape == CURVE_RIBBON) ? MTLCurveTypeFlat : MTLCurveTypeRound; - geomDescCrv.curveBasis = MTLCurveBasisCatmullRom; - geomDescCrv.curveEndCaps = MTLCurveEndCapsDisk; + if (hair->curve_shape == CURVE_THICK_LINEAR) { + geomDescCrv.curveBasis = MTLCurveBasisLinear; + geomDescCrv.curveEndCaps = MTLCurveEndCapsSphere; + } + else { + geomDescCrv.curveBasis = MTLCurveBasisCatmullRom; + geomDescCrv.curveEndCaps = MTLCurveEndCapsDisk; + } geomDescCrv.indexType = MTLIndexTypeUInt32; geomDescCrv.indexBuffer = idxBuffer; geomDescCrv.intersectionFunctionTableOffset = 1; diff --git a/intern/cycles/device/optix/device_impl.cpp b/intern/cycles/device/optix/device_impl.cpp index f5662c6778a..d468fc9d7f8 100644 --- a/intern/cycles/device/optix/device_impl.cpp +++ b/intern/cycles/device/optix/device_impl.cpp @@ -351,7 +351,8 @@ bool OptiXDevice::load_kernels(const uint kernel_features) 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_CATMULLROM; + pipeline_options.usesPrimitiveTypeFlags |= OPTIX_PRIMITIVE_TYPE_FLAGS_ROUND_LINEAR | + OPTIX_PRIMITIVE_TYPE_FLAGS_ROUND_CATMULLROM; } else { pipeline_options.usesPrimitiveTypeFlags |= OPTIX_PRIMITIVE_TYPE_FLAGS_CUSTOM; @@ -464,6 +465,29 @@ bool OptiXDevice::load_kernels(const uint kernel_features) group_descs[PG_HITS_MOTION] = group_descs[PG_HITS]; group_descs[PG_HITS_MOTION].hitgroup.moduleIS = builtin_modules[1]; } + + builtin_options.builtinISModuleType = OPTIX_PRIMITIVE_TYPE_ROUND_LINEAR; + builtin_options.usesMotionBlur = false; + + optix_assert(optixBuiltinISModuleGet( + context, &module_options, &pipeline_options, &builtin_options, &builtin_modules[2])); + + 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. */ @@ -613,6 +637,16 @@ bool OptiXDevice::load_kernels(const uint kernel_features) stack_size[PG_HITD_MOTION].cssIS + stack_size[PG_HITD_MOTION].cssAH); trace_css = std::max(trace_css, stack_size[PG_HITS_MOTION].cssIS + stack_size[PG_HITS_MOTION].cssAH); + trace_css = std::max( + trace_css, stack_size[PG_HITD_CURVE_LINEAR].cssIS + stack_size[PG_HITD_CURVE_LINEAR].cssAH); + trace_css = std::max( + trace_css, stack_size[PG_HITS_CURVE_LINEAR].cssIS + stack_size[PG_HITS_CURVE_LINEAR].cssAH); + trace_css = std::max(trace_css, + stack_size[PG_HITD_CURVE_LINEAR_MOTION].cssIS + + stack_size[PG_HITD_CURVE_LINEAR_MOTION].cssAH); + 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_POINTCLOUD].cssIS + stack_size[PG_HITD_POINTCLOUD].cssAH); trace_css = std::max( @@ -645,6 +679,14 @@ bool OptiXDevice::load_kernels(const uint kernel_features) pipeline_groups.push_back(groups[PG_HITD_MOTION]); pipeline_groups.push_back(groups[PG_HITS_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]); + if (pipeline_options.usesMotionBlur) { + pipeline_groups.push_back(groups[PG_HITD_CURVE_LINEAR_MOTION]); + pipeline_groups.push_back(groups[PG_HITS_CURVE_LINEAR_MOTION]); + } + } if (kernel_features & KERNEL_FEATURE_POINTCLOUD) { pipeline_groups.push_back(groups[PG_HITD_POINTCLOUD]); pipeline_groups.push_back(groups[PG_HITS_POINTCLOUD]); @@ -688,6 +730,14 @@ bool OptiXDevice::load_kernels(const uint kernel_features) pipeline_groups.push_back(groups[PG_HITD_MOTION]); pipeline_groups.push_back(groups[PG_HITS_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]); + if (pipeline_options.usesMotionBlur) { + pipeline_groups.push_back(groups[PG_HITD_CURVE_LINEAR_MOTION]); + pipeline_groups.push_back(groups[PG_HITS_CURVE_LINEAR_MOTION]); + } + } if (kernel_features & KERNEL_FEATURE_POINTCLOUD) { pipeline_groups.push_back(groups[PG_HITD_POINTCLOUD]); pipeline_groups.push_back(groups[PG_HITS_POINTCLOUD]); @@ -1221,7 +1271,12 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit) device_vector vertex_data(this, "optix temp vertex data", MEM_READ_ONLY); /* Four control points for each curve segment. */ size_t num_vertices = num_segments * 4; - if (hair->curve_shape == CURVE_THICK) { + if (hair->curve_shape == CURVE_THICK_LINEAR) { + num_vertices = hair->num_keys(); + index_data.alloc(num_segments); + vertex_data.alloc(num_vertices * num_motion_steps); + } + else if (hair->curve_shape == CURVE_THICK) { num_vertices = hair->num_keys() + 2 * hair->num_curves(); index_data.alloc(num_segments); vertex_data.alloc(num_vertices * num_motion_steps); @@ -1241,7 +1296,7 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit) keys = motion_keys->data_float3() + attr_offset * hair->get_curve_keys().size(); } - if (hair->curve_shape == CURVE_THICK) { + if (hair->curve_shape == CURVE_THICK || hair->curve_shape == CURVE_THICK_LINEAR) { for (size_t curve_index = 0, segment_index = 0, vertex_index = step * num_vertices; curve_index < hair->num_curves(); ++curve_index) @@ -1249,34 +1304,57 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit) const Hair::Curve curve = hair->get_curve(curve_index); const array &curve_radius = hair->get_curve_radius(); - const int first_key_index = curve.first_key; - { - vertex_data[vertex_index++] = make_float4(keys[first_key_index].x, - keys[first_key_index].y, - keys[first_key_index].z, - curve_radius[first_key_index]); - } + if (hair->curve_shape == CURVE_THICK_LINEAR) { + const int first_key_index = curve.first_key; - for (int k = 0; k < curve.num_segments(); ++k) { - if (step == 0) { - index_data[segment_index++] = vertex_index - 1; + for (int k = 0; k < curve.num_segments(); ++k) { + if (step == 0) { + index_data[segment_index++] = vertex_index; + } + vertex_data[vertex_index++] = make_float4(keys[first_key_index + k].x, + keys[first_key_index + k].y, + keys[first_key_index + k].z, + curve_radius[first_key_index + k]); } - vertex_data[vertex_index++] = make_float4(keys[first_key_index + k].x, - keys[first_key_index + k].y, - keys[first_key_index + k].z, - curve_radius[first_key_index + k]); - } - const int last_key_index = first_key_index + curve.num_keys - 1; - { - vertex_data[vertex_index++] = make_float4(keys[last_key_index].x, - keys[last_key_index].y, - keys[last_key_index].z, - curve_radius[last_key_index]); - vertex_data[vertex_index++] = make_float4(keys[last_key_index].x, - keys[last_key_index].y, - keys[last_key_index].z, - curve_radius[last_key_index]); + const int last_key_index = first_key_index + curve.num_keys - 1; + { + vertex_data[vertex_index++] = make_float4(keys[last_key_index].x, + keys[last_key_index].y, + keys[last_key_index].z, + curve_radius[last_key_index]); + } + } + else { + const int first_key_index = curve.first_key; + { + vertex_data[vertex_index++] = make_float4(keys[first_key_index].x, + keys[first_key_index].y, + keys[first_key_index].z, + curve_radius[first_key_index]); + } + + for (int k = 0; k < curve.num_segments(); ++k) { + if (step == 0) { + index_data[segment_index++] = vertex_index - 1; + } + vertex_data[vertex_index++] = make_float4(keys[first_key_index + k].x, + keys[first_key_index + k].y, + keys[first_key_index + k].z, + curve_radius[first_key_index + k]); + } + + const int last_key_index = first_key_index + curve.num_keys - 1; + { + vertex_data[vertex_index++] = make_float4(keys[last_key_index].x, + keys[last_key_index].y, + keys[last_key_index].z, + curve_radius[last_key_index]); + vertex_data[vertex_index++] = make_float4(keys[last_key_index].x, + keys[last_key_index].y, + keys[last_key_index].z, + curve_radius[last_key_index]); + } } } } @@ -1322,9 +1400,14 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit) /* Force a single any-hit call, so shadow record-all behavior works correctly. */ unsigned int build_flags = OPTIX_GEOMETRY_FLAG_REQUIRE_SINGLE_ANYHIT_CALL; OptixBuildInput build_input = {}; - if (hair->curve_shape == CURVE_THICK) { + if (hair->curve_shape != CURVE_RIBBON) { build_input.type = OPTIX_BUILD_INPUT_TYPE_CURVES; - build_input.curveArray.curveType = OPTIX_PRIMITIVE_TYPE_ROUND_CATMULLROM; + if (hair->curve_shape == CURVE_THICK_LINEAR) { + build_input.curveArray.curveType = OPTIX_PRIMITIVE_TYPE_ROUND_LINEAR; + } + else { + build_input.curveArray.curveType = OPTIX_PRIMITIVE_TYPE_ROUND_CATMULLROM; + } build_input.curveArray.numPrimitives = num_segments; build_input.curveArray.vertexBuffers = (CUdeviceptr *)vertex_ptrs.data(); build_input.curveArray.numVertices = num_vertices; @@ -1589,11 +1672,19 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit) } if (ob->get_geometry()->is_hair() && - static_cast(ob->get_geometry())->curve_shape == CURVE_THICK) + static_cast(ob->get_geometry())->curve_shape != CURVE_RIBBON) { - 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; + if (static_cast(ob->get_geometry())->curve_shape == CURVE_THICK_LINEAR) { + instance.sbtOffset = PG_HITD_CURVE_LINEAR - PG_HITD; + if (pipeline_options.usesMotionBlur && ob->get_geometry()->has_motion_blur()) { + instance.sbtOffset = PG_HITD_CURVE_LINEAR_MOTION - PG_HITD; + } + } + else { + 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; + } } } else if (ob->get_geometry()->is_pointcloud()) { diff --git a/intern/cycles/device/optix/device_impl.h b/intern/cycles/device/optix/device_impl.h index d2d6d069124..a1570310e28 100644 --- a/intern/cycles/device/optix/device_impl.h +++ b/intern/cycles/device/optix/device_impl.h @@ -44,6 +44,10 @@ enum { PG_HITV, /* __VOLUME__ hit group. */ PG_HITD_MOTION, PG_HITS_MOTION, + PG_HITD_CURVE_LINEAR, + PG_HITS_CURVE_LINEAR, + PG_HITD_CURVE_LINEAR_MOTION, + PG_HITS_CURVE_LINEAR_MOTION, PG_HITD_POINTCLOUD, PG_HITS_POINTCLOUD, PG_CALL_SVM_AO, @@ -71,7 +75,7 @@ class OptiXDevice : public CUDADevice { OptixDeviceContext context = nullptr; OptixModule optix_module = nullptr; /* All necessary OptiX kernels are in one module. */ - OptixModule builtin_modules[2] = {}; + OptixModule builtin_modules[4] = {}; OptixPipeline pipelines[NUM_PIPELINES] = {}; OptixProgramGroup groups[NUM_PROGRAM_GROUPS] = {}; OptixPipelineCompileOptions pipeline_options = {}; diff --git a/intern/cycles/kernel/bvh/shadow_all.h b/intern/cycles/kernel/bvh/shadow_all.h index 1f020a24340..d56d3c363ed 100644 --- a/intern/cycles/kernel/bvh/shadow_all.h +++ b/intern/cycles/kernel/bvh/shadow_all.h @@ -181,7 +181,9 @@ ccl_device_inline case PRIMITIVE_CURVE_THICK: case PRIMITIVE_MOTION_CURVE_THICK: case PRIMITIVE_CURVE_RIBBON: - case PRIMITIVE_MOTION_CURVE_RIBBON: { + case PRIMITIVE_MOTION_CURVE_RIBBON: + case PRIMITIVE_CURVE_THICK_LINEAR: + case PRIMITIVE_MOTION_CURVE_THICK_LINEAR: { if ((type & PRIMITIVE_MOTION) && kernel_data.bvh.use_bvh_steps) { const float2 prim_time = kernel_data_fetch(prim_time, prim_addr); if (ray->time < prim_time.x || ray->time > prim_time.y) { diff --git a/intern/cycles/kernel/bvh/traversal.h b/intern/cycles/kernel/bvh/traversal.h index dc9829be55a..86030691a14 100644 --- a/intern/cycles/kernel/bvh/traversal.h +++ b/intern/cycles/kernel/bvh/traversal.h @@ -182,7 +182,9 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg, case PRIMITIVE_CURVE_THICK: case PRIMITIVE_MOTION_CURVE_THICK: case PRIMITIVE_CURVE_RIBBON: - case PRIMITIVE_MOTION_CURVE_RIBBON: { + case PRIMITIVE_MOTION_CURVE_RIBBON: + case PRIMITIVE_CURVE_THICK_LINEAR: + case PRIMITIVE_MOTION_CURVE_THICK_LINEAR: { if ((type & PRIMITIVE_MOTION) && kernel_data.bvh.use_bvh_steps) { const float2 prim_time = kernel_data_fetch(prim_time, prim_addr); if (ray->time < prim_time.x || ray->time > prim_time.y) { diff --git a/intern/cycles/kernel/closure/bsdf_principled_hair_chiang.h b/intern/cycles/kernel/closure/bsdf_principled_hair_chiang.h index 4ac947faca6..d27480e2b9d 100644 --- a/intern/cycles/kernel/closure/bsdf_principled_hair_chiang.h +++ b/intern/cycles/kernel/closure/bsdf_principled_hair_chiang.h @@ -181,7 +181,8 @@ ccl_device int bsdf_hair_chiang_setup(ccl_private ShaderData *sd, ccl_private Ch /* TODO: we convert this value to a cosine later and discard the sign, so * we could probably save some operations. */ - bsdf->h = (sd->type & PRIMITIVE_CURVE_RIBBON) ? -sd->v : dot(cross(sd->Ng, X), Z); + bsdf->h = ((sd->type & PRIMITIVE_CURVE) == PRIMITIVE_CURVE_RIBBON) ? -sd->v : + dot(cross(sd->Ng, X), Z); kernel_assert(fabsf(bsdf->h) < 1.0f + 1e-4f); kernel_assert(isfinite_safe(Y)); diff --git a/intern/cycles/kernel/closure/bsdf_principled_hair_huang.h b/intern/cycles/kernel/closure/bsdf_principled_hair_huang.h index f275c8cf4be..cb235a53908 100644 --- a/intern/cycles/kernel/closure/bsdf_principled_hair_huang.h +++ b/intern/cycles/kernel/closure/bsdf_principled_hair_huang.h @@ -257,7 +257,7 @@ ccl_device int bsdf_hair_huang_setup(ccl_private ShaderData *sd, /* h from -1..0..1 means the rays goes from grazing the hair, to hitting it at the center, to * grazing the other edge. This is the cosine of the angle between `sd->N` and `X`. */ - bsdf->h = (sd->type & PRIMITIVE_CURVE_RIBBON) ? -sd->v : -dot(X, sd->N); + bsdf->h = ((sd->type & PRIMITIVE_CURVE) == PRIMITIVE_CURVE_RIBBON) ? -sd->v : -dot(X, sd->N); kernel_assert(fabsf(bsdf->h) < 1.0f + 1e-4f); kernel_assert(isfinite_safe(bsdf->h)); diff --git a/intern/cycles/kernel/device/cpu/bvh.h b/intern/cycles/kernel/device/cpu/bvh.h index cf01b316f18..ca043ef5e09 100644 --- a/intern/cycles/kernel/device/cpu/bvh.h +++ b/intern/cycles/kernel/device/cpu/bvh.h @@ -51,7 +51,8 @@ using numhit_t = uint32_t; (RTCFeatureFlags)(RTC_FEATURE_FLAG_TRIANGLE | RTC_FEATURE_FLAG_INSTANCE | \ RTC_FEATURE_FLAG_FILTER_FUNCTION_IN_ARGUMENTS | RTC_FEATURE_FLAG_POINT | \ RTC_FEATURE_FLAG_MOTION_BLUR | RTC_FEATURE_FLAG_ROUND_CATMULL_ROM_CURVE | \ - RTC_FEATURE_FLAG_FLAT_CATMULL_ROM_CURVE) + RTC_FEATURE_FLAG_FLAT_CATMULL_ROM_CURVE | \ + RTC_FEATURE_FLAG_ROUND_LINEAR_CURVE) #endif #define EMBREE_IS_HAIR(x) (x & 1) diff --git a/intern/cycles/kernel/device/metal/kernel.metal b/intern/cycles/kernel/device/metal/kernel.metal index 28182f50ebe..642c9bfd208 100644 --- a/intern/cycles/kernel/device/metal/kernel.metal +++ b/intern/cycles/kernel/device/metal/kernel.metal @@ -176,6 +176,11 @@ __intersection__local_tri_mblur( payload, primitive_id, barycentrics, ray_tmax); } +inline bool metalrt_curve_skip_end_cap(const int type, const float u) +{ + return ((u == 0.0f || u == 1.0f) && (type & PRIMITIVE_CURVE) != PRIMITIVE_CURVE_THICK_LINEAR); +} + template bool metalrt_shadow_all_hit( constant KernelParamsMetal &launch_params_metal, @@ -205,7 +210,7 @@ bool metalrt_shadow_all_hit( prim = segment.prim; /* Filter out curve end-caps. */ - if (u == 0.0f || u == 1.0f) { + if (metalrt_curve_skip_end_cap(type, u)) { /* continue search */ return true; } @@ -405,17 +410,17 @@ inline TReturnType metalrt_visibility_test( # ifdef __HAIR__ if constexpr (intersection_type == METALRT_HIT_CURVE) { + const KernelCurveSegment segment = kernel_data_fetch(curve_segments, prim); + int type = segment.type; + prim = segment.prim; + /* Filter out curve end-caps. */ - if (u == 0.0f || u == 1.0f) { + if (metalrt_curve_skip_end_cap(type, u)) { result.accept = false; result.continue_search = true; return result; } - const KernelCurveSegment segment = kernel_data_fetch(curve_segments, prim); - int type = segment.type; - prim = segment.prim; - if (type & PRIMITIVE_CURVE_RIBBON) { MetalKernelContext context(launch_params_metal); if (!context.curve_ribbon_accept(nullptr, u, t, ray, object, prim, type)) { @@ -456,17 +461,17 @@ inline TReturnType metalrt_visibility_test_shadow( # ifdef __HAIR__ if constexpr (intersection_type == METALRT_HIT_CURVE) { + const KernelCurveSegment segment = kernel_data_fetch(curve_segments, prim); + int type = segment.type; + prim = segment.prim; + /* Filter out curve end-caps. */ - if (u == 0.0f || u == 1.0f) { + if (metalrt_curve_skip_end_cap(type, u)) { result.accept = false; result.continue_search = true; return result; } - const KernelCurveSegment segment = kernel_data_fetch(curve_segments, prim); - int type = segment.type; - prim = segment.prim; - if (type & PRIMITIVE_CURVE_RIBBON) { MetalKernelContext context(launch_params_metal); if (!context.curve_ribbon_accept(nullptr, u, t, ray, object, prim, type)) { diff --git a/intern/cycles/kernel/device/oneapi/kernel.cpp b/intern/cycles/kernel/device/oneapi/kernel.cpp index bcd53bc18fd..8f09961267b 100644 --- a/intern/cycles/kernel/device/oneapi/kernel.cpp +++ b/intern/cycles/kernel/device/oneapi/kernel.cpp @@ -31,7 +31,7 @@ static const RTCFeatureFlags CYCLES_ONEAPI_EMBREE_BASIC_FEATURES = (const RTCFea RTC_FEATURE_FLAG_MOTION_BLUR); static const RTCFeatureFlags CYCLES_ONEAPI_EMBREE_ALL_FEATURES = (const RTCFeatureFlags)( CYCLES_ONEAPI_EMBREE_BASIC_FEATURES | RTC_FEATURE_FLAG_ROUND_CATMULL_ROM_CURVE | - RTC_FEATURE_FLAG_FLAT_CATMULL_ROM_CURVE); + RTC_FEATURE_FLAG_FLAT_CATMULL_ROM_CURVE | RTC_FEATURE_FLAG_ROUND_LINEAR_CURVE); # endif void oneapi_set_error_cb(OneAPIErrorCallback cb, void *user_ptr) diff --git a/intern/cycles/kernel/device/optix/bvh.h b/intern/cycles/kernel/device/optix/bvh.h index b301233a8f1..de253ca5dd3 100644 --- a/intern/cycles/kernel/device/optix/bvh.h +++ b/intern/cycles/kernel/device/optix/bvh.h @@ -163,13 +163,6 @@ extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit() const KernelCurveSegment segment = kernel_data_fetch(curve_segments, prim); type = segment.type; prim = segment.prim; - -# if OPTIX_ABI_VERSION < 55 - /* Filter out curve end-caps. */ - if (u == 0.0f || u == 1.0f) { - return optixIgnoreIntersection(); - } -# endif } # endif else { @@ -306,18 +299,6 @@ extern "C" __global__ void __anyhit__kernel_optix_volume_test() extern "C" __global__ void __anyhit__kernel_optix_visibility_test() { -#ifdef __HAIR__ -# if OPTIX_ABI_VERSION < 55 - if (optixGetPrimitiveType() == OPTIX_PRIMITIVE_TYPE_ROUND_CUBIC_BSPLINE) { - /* Filter out curve end-caps. */ - const float u = __uint_as_float(optixGetAttribute_0()); - if (u == 0.0f || u == 1.0f) { - return optixIgnoreIntersection(); - } - } -# endif -#endif - const uint object = get_object_id(); const uint visibility = optixGetPayload_4(); #ifdef __VISIBILITY_FLAG__ diff --git a/intern/cycles/kernel/geom/curve_intersect.h b/intern/cycles/kernel/geom/curve_intersect.h index fd61dd008a0..f3b24ba97d2 100644 --- a/intern/cycles/kernel/geom/curve_intersect.h +++ b/intern/cycles/kernel/geom/curve_intersect.h @@ -1,4 +1,4 @@ -/* SPDX-FileCopyrightText: 2009-2020 Intel Corporation +/* SPDX-FileCopyrightText: 2009-2021 Intel Corporation * * SPDX-License-Identifier: Apache-2.0 * @@ -51,7 +51,6 @@ ccl_device_inline float4 catmull_rom_basis_derivative(const float4 curve[4], flo ccl_device_inline float4 catmull_rom_basis_derivative2(const float4 curve[4], float u) { - const float t = u; const float n0 = -3.0f * t + 2.0f; const float n1 = 9.0f * t - 5.0f; @@ -634,6 +633,168 @@ ccl_device_inline bool ribbon_intersect(const float3 ray_org, return false; } +/* Linear curve evaluation. */ + +ccl_device_inline float4 linear_basis_eval(const float4 curve[4], float u) +{ + return mix(curve[1], curve[2], u); +} + +ccl_device_inline float4 linear_basis_derivative(const float4 curve[4], float) +{ + return curve[2] - curve[1]; +} + +/* Linear Thick Curve */ + +ccl_device_inline bool cone_sphere_intersect(const float4 curve[4], + const float3 ray_D, + ccl_private float *t_o, + ccl_private float *u_o, + ccl_private float3 *Ng_o) +{ + /* Calculate quadratic equation to solve. */ + const float r0 = curve[1].w; + const float r1 = curve[2].w; + const float dr = r1 - r0; + const float r0dr = r0 * dr; + + const float3 P0 = make_float3(curve[1]); + const float3 P1 = make_float3(curve[2]); + const float3 dP = P1 - P0; + const float3 O = -P0; + const float3 dO = ray_D; + + const float dOdO = dot(dO, dO); + const float OdO = dot(dO, O); + const float OO = dot(O, O); + const float dOz = dot(dP, dO); + const float Oz = dot(dP, O); + + const float dPdP = dot(dP, dP); + const float yp = Oz + r0dr; + const float g = dPdP - sqr(dr); + + const float A = g * dOdO - sqr(dOz); + const float B = 2.0f * (g * OdO - dOz * yp); + const float C = g * OO - sqr(Oz) - sqr(r0) * dPdP - 2.0f * r0dr * Oz; + + /* We miss the cone if determinant is smaller than zero. */ + const float D = B * B - 4.0f * A * C; + if (!(D >= 0.0f)) { + *t_o = FLT_MAX; + return false; + } + + /* Special case for rays that are parallel to the cone. */ + const float eps = 1e-18f; + if (fabsf(A) < eps) { + *t_o = -FLT_MAX; + return false; + } + + /* Standard case for rays that are not parallel to the cone. */ + const float Q = sqrtf(D); + const float rcp_2A = 1.0f / (2.0f * A); + const float t0 = (-B - Q) * rcp_2A; + const float y0 = yp + t0 * dOz; + + float t = FLT_MAX; + + /* Calculates u and Ng for near hit. */ + if ((y0 > -FLT_EPSILON) && (y0 <= g) && (g > 0.0f)) { + t = t0; + *u_o = clamp(y0 / g, 0.0f, 1.0f); + const float3 Pr = O + t0 * dO; + const float3 Pl = (*u_o) * dP; + *Ng_o = Pr - Pl; + } + + /* Intersect ending sphere. */ + { + const float3 O1 = -P1; + const float O1dO = dot(O1, dO); + const float h2 = sqr(O1dO) - dOdO * (dot(O1, O1) - sqr(r1)); + if (h2 >= 0.0f) { + const float rhs1 = sqrt(h2); + + /* Clip away near hit if it is inside next cone segment. */ + const float t_sph1 = (-O1dO - rhs1) * (1.0f / dOdO); + + const float r2 = curve[3].w; + const float3 P2 = make_float3(curve[3]); + const float y2 = dot((t_sph1 * dO) - P1, (P2 - P1)); + const float cap2 = -(r1 * (r2 - r1)); + + if ((t_sph1 <= t) && (yp + t_sph1 * dOz) > g && !(y2 > cap2)) { + t = t_sph1; + *u_o = 1.0f; + *Ng_o = t * dO - P1; + } + } + } + + /* Intersect start sphere. */ + if (isequal(curve[0], curve[1])) { + const float h2 = sqr(OdO) - dOdO * (dot(O, O) - sqr(r0)); + if (h2 >= 0.0f) { + const float rhs1 = sqrt(h2); + + /* Clip away near hit if it is inside next cone segment. */ + const float t_sph0 = (-OdO - rhs1) * (1.0f / dOdO); + + if ((t_sph0 <= t) && (yp + t_sph0 * dOz) < 0) { + t = t_sph0; + *u_o = 0.0f; + *Ng_o = t * dO - P0; + } + } + } + + *t_o = t; + + return t != FLT_MAX; +} + +ccl_device bool linear_curve_intersect(const float3 ray_P, + const float3 ray_D, + const float ray_tmin, + float ray_tmax, + float4 curve[4], + ccl_private Intersection *isect) +{ + /* Move ray closer to make intersection stable. */ + const float3 center = make_float3(0.5f * (curve[1] + curve[2])); + const float dt = dot(center - ray_P, ray_D) / dot(ray_D, ray_D); + const float3 ref = ray_P + ray_D * dt; + const float4 ref4 = make_float4(ref, 0.0f); + curve[0] -= ref4; + curve[1] -= ref4; + curve[2] -= ref4; + curve[3] -= ref4; + + /* Intersect with cone sphere. */ + float t; + float u; + float3 Ng; + if (!cone_sphere_intersect(curve, ray_D, &t, &u, &Ng)) { + return false; + } + + t += dt; + + if (!(t >= ray_tmin && t <= ray_tmax)) { + return false; /* Rejects NaNs */ + } + + /* Record intersection. */ + isect->t = t; + isect->u = u; + isect->v = 0.0f; + + return true; +} + ccl_device_forceinline bool curve_intersect(KernelGlobals kg, ccl_private Intersection *isect, const float3 ray_P, @@ -665,24 +826,36 @@ ccl_device_forceinline bool curve_intersect(KernelGlobals kg, motion_curve_keys(kg, object, time, ka, k0, k1, kb, curve); } - if (type & PRIMITIVE_CURVE_RIBBON) { - /* todo: adaptive number of subdivisions could help performance here. */ - const int subdivisions = kernel_data.bvh.curve_subdivisions; - if (ribbon_intersect(ray_P, ray_D, tmin, tmax, subdivisions, curve, isect)) { - isect->prim = prim; - isect->object = object; - isect->type = type; - return true; + switch (type & PRIMITIVE_CURVE) { + case PRIMITIVE_CURVE_RIBBON: { + /* todo: adaptive number of subdivisions could help performance here. */ + const int subdivisions = kernel_data.bvh.curve_subdivisions; + if (ribbon_intersect(ray_P, ray_D, tmin, tmax, subdivisions, curve, isect)) { + isect->prim = prim; + isect->object = object; + isect->type = type; + return true; + } + break; + } + case PRIMITIVE_CURVE_THICK: { + if (curve_intersect_recursive(ray_P, ray_D, tmin, tmax, curve, isect)) { + isect->prim = prim; + isect->object = object; + isect->type = type; + return true; + } + break; + } + case PRIMITIVE_CURVE_THICK_LINEAR: { + if (linear_curve_intersect(ray_P, ray_D, tmin, tmax, curve, isect)) { + isect->prim = prim; + isect->object = object; + isect->type = type; + return true; + } + break; } - - return false; - } - - if (curve_intersect_recursive(ray_P, ray_D, tmin, tmax, curve, isect)) { - isect->prim = prim; - isect->object = object; - isect->type = type; - return true; } return false; @@ -724,10 +897,12 @@ ccl_device_inline void curve_shader_setup(KernelGlobals kg, P = P + D * t; - const float4 dPdu4 = catmull_rom_basis_derivative(P_curve, sd->u); + const float4 dPdu4 = (sd->type & PRIMITIVE_CURVE) == PRIMITIVE_CURVE_THICK_LINEAR ? + linear_basis_derivative(P_curve, sd->u) : + catmull_rom_basis_derivative(P_curve, sd->u); const float3 dPdu = make_float3(dPdu4); - if (sd->type & PRIMITIVE_CURVE_RIBBON) { + if ((sd->type & PRIMITIVE_CURVE) == PRIMITIVE_CURVE_RIBBON) { /* Rounded smooth normals for ribbons, to approximate thick curve shape. */ const float3 tangent = normalize(dPdu); const float3 bitangent = normalize(cross(tangent, -D)); @@ -749,7 +924,9 @@ ccl_device_inline void curve_shader_setup(KernelGlobals kg, * however for Optix this would go beyond the size of the payload. */ /* NOTE: It is possible that P will be the same as P_inside (precision issues, or very small * radius). In this case use the view direction to approximate the normal. */ - const float3 P_inside = make_float3(catmull_rom_basis_eval(P_curve, sd->u)); + const float3 P_inside = (sd->type & PRIMITIVE_CURVE) == PRIMITIVE_CURVE_THICK_LINEAR ? + make_float3(linear_basis_eval(P_curve, sd->u)) : + make_float3(catmull_rom_basis_eval(P_curve, sd->u)); const float3 N = (!isequal(P, P_inside)) ? normalize(P - P_inside) : -sd->wi; sd->N = N; @@ -769,7 +946,7 @@ ccl_device_inline void curve_shader_setup(KernelGlobals kg, } sd->P = P; - sd->Ng = (sd->type & PRIMITIVE_CURVE_RIBBON) ? sd->wi : sd->N; + sd->Ng = ((sd->type & PRIMITIVE_CURVE) == PRIMITIVE_CURVE_RIBBON) ? sd->wi : sd->N; sd->dPdv = cross(sd->dPdu, sd->Ng); sd->shader = kernel_data_fetch(curves, sd->prim).shader_id; } diff --git a/intern/cycles/kernel/geom/shader_data.h b/intern/cycles/kernel/geom/shader_data.h index b8e6d449238..e2eb6727fb3 100644 --- a/intern/cycles/kernel/geom/shader_data.h +++ b/intern/cycles/kernel/geom/shader_data.h @@ -325,9 +325,13 @@ ccl_device void shader_setup_from_curve(KernelGlobals kg, P_curve[3] = kernel_data_fetch(curve_keys, kb); /* Interpolate position and tangent. */ - sd->P = make_float3(catmull_rom_basis_eval(P_curve, sd->u)); + sd->P = (sd->type & PRIMITIVE_CURVE) == PRIMITIVE_CURVE_THICK_LINEAR ? + make_float3(linear_basis_eval(P_curve, sd->u)) : + make_float3(catmull_rom_basis_eval(P_curve, sd->u)); # ifdef __DPDU__ - sd->dPdu = make_float3(catmull_rom_basis_derivative(P_curve, sd->u)); + sd->dPdu = (sd->type & PRIMITIVE_CURVE) == PRIMITIVE_CURVE_THICK_LINEAR ? + make_float3(linear_basis_derivative(P_curve, sd->u)) : + make_float3(catmull_rom_basis_derivative(P_curve, sd->u)); # endif /* Transform into world space */ diff --git a/intern/cycles/kernel/types.h b/intern/cycles/kernel/types.h index ce47b3af627..ac85c0efc3d 100644 --- a/intern/cycles/kernel/types.h +++ b/intern/cycles/kernel/types.h @@ -811,6 +811,7 @@ enum PrimitiveType { PRIMITIVE_TRIANGLE = (1 << 0), PRIMITIVE_CURVE_THICK = (1 << 1), PRIMITIVE_CURVE_RIBBON = (1 << 2), + PRIMITIVE_CURVE_THICK_LINEAR = PRIMITIVE_CURVE_THICK | PRIMITIVE_CURVE_RIBBON, PRIMITIVE_POINT = (1 << 3), PRIMITIVE_VOLUME = (1 << 4), PRIMITIVE_LAMP = (1 << 5), @@ -819,6 +820,7 @@ enum PrimitiveType { PRIMITIVE_MOTION_TRIANGLE = (PRIMITIVE_TRIANGLE | PRIMITIVE_MOTION), PRIMITIVE_MOTION_CURVE_THICK = (PRIMITIVE_CURVE_THICK | PRIMITIVE_MOTION), PRIMITIVE_MOTION_CURVE_RIBBON = (PRIMITIVE_CURVE_RIBBON | PRIMITIVE_MOTION), + PRIMITIVE_MOTION_CURVE_THICK_LINEAR = (PRIMITIVE_CURVE_THICK_LINEAR | PRIMITIVE_MOTION), PRIMITIVE_MOTION_POINT = (PRIMITIVE_POINT | PRIMITIVE_MOTION), PRIMITIVE_CURVE = (PRIMITIVE_CURVE_THICK | PRIMITIVE_CURVE_RIBBON), @@ -842,6 +844,7 @@ enum PrimitiveType { enum CurveShapeType { CURVE_RIBBON = 0, CURVE_THICK = 1, + CURVE_THICK_LINEAR = 2, CURVE_NUM_SHAPE_TYPES, }; diff --git a/intern/cycles/scene/geometry.cpp b/intern/cycles/scene/geometry.cpp index 3ac223bebc5..20dc96fa98a 100644 --- a/intern/cycles/scene/geometry.cpp +++ b/intern/cycles/scene/geometry.cpp @@ -432,7 +432,9 @@ 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); - hair->curve_shape = scene->params.hair_shape; + 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/hair.cpp b/intern/cycles/scene/hair.cpp index d3d3aaa53b0..10bae90b6bf 100644 --- a/intern/cycles/scene/hair.cpp +++ b/intern/cycles/scene/hair.cpp @@ -529,9 +529,12 @@ void Hair::pack_curves(Scene *scene, PrimitiveType Hair::primitive_type() const { return has_motion_blur() ? - ((curve_shape == CURVE_RIBBON) ? PRIMITIVE_MOTION_CURVE_RIBBON : - PRIMITIVE_MOTION_CURVE_THICK) : - ((curve_shape == CURVE_RIBBON) ? PRIMITIVE_CURVE_RIBBON : PRIMITIVE_CURVE_THICK); + ((curve_shape == CURVE_RIBBON) ? PRIMITIVE_MOTION_CURVE_RIBBON : + (curve_shape == CURVE_THICK_LINEAR) ? PRIMITIVE_MOTION_CURVE_THICK_LINEAR : + PRIMITIVE_MOTION_CURVE_THICK) : + ((curve_shape == CURVE_RIBBON) ? PRIMITIVE_CURVE_RIBBON : + (curve_shape == CURVE_THICK_LINEAR) ? PRIMITIVE_CURVE_THICK_LINEAR : + PRIMITIVE_CURVE_THICK); } /* Fill in coordinates for curve transparency shader evaluation on device. */ diff --git a/intern/cycles/scene/scene.cpp b/intern/cycles/scene/scene.cpp index ca6b1de444e..bd7ea76a8d4 100644 --- a/intern/cycles/scene/scene.cpp +++ b/intern/cycles/scene/scene.cpp @@ -500,7 +500,7 @@ 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) { + if (params.hair_shape == CURVE_THICK || params.hair_shape == CURVE_THICK_LINEAR) { kernel_features |= KERNEL_FEATURE_HAIR_THICK; } diff --git a/tests/files/render/hair/cycles_renders/hair_linear_close_up.png b/tests/files/render/hair/cycles_renders/hair_linear_close_up.png new file mode 100644 index 00000000000..4934b4cb6e8 --- /dev/null +++ b/tests/files/render/hair/cycles_renders/hair_linear_close_up.png @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:f980b4bf0e9780d67d6bb7af384c87daa9b7186754f3f52c5c58d18386c104dd +size 14775 diff --git a/tests/files/render/hair/eevee_renders/hair_linear_close_up.png b/tests/files/render/hair/eevee_renders/hair_linear_close_up.png new file mode 100644 index 00000000000..fb7519d557c --- /dev/null +++ b/tests/files/render/hair/eevee_renders/hair_linear_close_up.png @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:c513bdd910ea8b13ab345675265c725730a691b7c98092ecc1f472351277ad03 +size 15024 diff --git a/tests/files/render/hair/hair_linear_close_up.blend b/tests/files/render/hair/hair_linear_close_up.blend new file mode 100644 index 00000000000..55ba216c9f9 --- /dev/null +++ b/tests/files/render/hair/hair_linear_close_up.blend @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:1ad67e04d7dc1d81e8b6529be487cb3121d53775c241ccc5be55908224b55382 +size 145964 diff --git a/tests/files/render/hair/storm_hydra_renders/hair_linear_close_up.png b/tests/files/render/hair/storm_hydra_renders/hair_linear_close_up.png new file mode 100644 index 00000000000..5d6db73f88b --- /dev/null +++ b/tests/files/render/hair/storm_hydra_renders/hair_linear_close_up.png @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:84eb0a305ff8735f5b85d70ac240fed6baab44372bff2846c9cdf5ac93e73200 +size 1672 diff --git a/tests/files/render/hair/storm_usd_renders/hair_linear_close_up.png b/tests/files/render/hair/storm_usd_renders/hair_linear_close_up.png new file mode 100644 index 00000000000..d31b97c0e09 --- /dev/null +++ b/tests/files/render/hair/storm_usd_renders/hair_linear_close_up.png @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:a893f9b4f70c24a568a288d8151bacd545c4bb72b9fd1afecf21e0197b372001 +size 885 diff --git a/tests/files/render/hair/workbench_renders/hair_linear_close_up.png b/tests/files/render/hair/workbench_renders/hair_linear_close_up.png new file mode 100644 index 00000000000..a9a87a9fd9e --- /dev/null +++ b/tests/files/render/hair/workbench_renders/hair_linear_close_up.png @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:061b4ae88786d7c268954c1598b35995a743ca923f68397f6daf936726f05621 +size 12464