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

Cycles: Add linear curve shape

Browse Source 
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
This commit is contained in:
Patrick Mours
2025-07-29 17:05:01 +02:00
committed by Brecht Van Lommel
parent 0a1d146389
commit 6487395fa5
26 changed files with 480 additions and 133 deletions
Download Patch File Download Diff File Expand all files Collapse all files

11
intern/cycles/blender/addon/properties.py
View File

@@ -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 = (

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

@@ -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<int8_t> 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

50
intern/cycles/bvh/embree.cpp
View File

@@ -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<float3>(num_curves, hair, verts, curve_radius, rtc_verts);
pack_motion_verts<float3>(
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<float4>(num_curves, hair, verts, curve_radius, rtc_verts);
pack_motion_verts<float4>(
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;
}

48
intern/cycles/device/metal/bvh.mm
View File

@@ -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;

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

@@ -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<float4> 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<float> &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<const Hair *>(ob->get_geometry())->curve_shape == CURVE_THICK)
static_cast<const Hair *>(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<const Hair *>(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()) {

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

@@ -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 = {};

4
intern/cycles/kernel/bvh/shadow_all.h
View File

@@ -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) {

4
intern/cycles/kernel/bvh/traversal.h
View File

@@ -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) {

3
intern/cycles/kernel/closure/bsdf_principled_hair_chiang.h
View File

@@ -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));

2
intern/cycles/kernel/closure/bsdf_principled_hair_huang.h
View File

@@ -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));

3
intern/cycles/kernel/device/cpu/bvh.h
View File

@@ -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)

27
intern/cycles/kernel/device/metal/kernel.metal
View File

@@ -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<uint intersection_type>
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)) {

2
intern/cycles/kernel/device/oneapi/kernel.cpp
View File

@@ -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)

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

@@ -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__

223
intern/cycles/kernel/geom/curve_intersect.h
View File

@@ -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;
}

8
intern/cycles/kernel/geom/shader_data.h
View File

@@ -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 */

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

@@ -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,
};

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

@@ -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<Hair *>(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;

9
intern/cycles/scene/hair.cpp
View File

@@ -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. */

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

@@ -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;
}

BIN
tests/files/render/hair/cycles_renders/hair_linear_close_up.png (Stored with Git LFS) Normal file
View File

Binary file not shown.

BIN
tests/files/render/hair/eevee_renders/hair_linear_close_up.png (Stored with Git LFS) Normal file
View File

Binary file not shown.

BIN
tests/files/render/hair/hair_linear_close_up.blend (Stored with Git LFS) Normal file
View File

Binary file not shown.

BIN
tests/files/render/hair/storm_hydra_renders/hair_linear_close_up.png (Stored with Git LFS) Normal file
View File

Binary file not shown.

BIN
tests/files/render/hair/storm_usd_renders/hair_linear_close_up.png (Stored with Git LFS) Normal file
View File

Binary file not shown.

BIN
tests/files/render/hair/workbench_renders/hair_linear_close_up.png (Stored with Git LFS) Normal file
View File

Binary file not shown.