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