Welcome to mirror list, hosted at ThFree Co, Russian Federation.

git.blender.org/blender.git - Unnamed repository; edit this file 'description' to name the repository.
summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
Diffstat (limited to 'intern/cycles/kernel/kernels/optix/kernel_optix.cu')
-rw-r--r--intern/cycles/kernel/kernels/optix/kernel_optix.cu22
1 files changed, 19 insertions, 3 deletions
diff --git a/intern/cycles/kernel/kernels/optix/kernel_optix.cu b/intern/cycles/kernel/kernels/optix/kernel_optix.cu
index 59fece73b3c..c730d952ed4 100644
--- a/intern/cycles/kernel/kernels/optix/kernel_optix.cu
+++ b/intern/cycles/kernel/kernels/optix/kernel_optix.cu
@@ -256,11 +256,9 @@ extern "C" __global__ void __closesthit__kernel_optix_hit()
}
#ifdef __HAIR__
-extern "C" __global__ void __intersection__curve()
+ccl_device_inline void optix_intersection_curve(const uint prim, const uint type)
{
- const uint prim = optixGetPrimitiveIndex();
const uint object = get_object_id<true>();
- const uint type = kernel_tex_fetch(__prim_type, prim);
const uint visibility = optixGetPayload_4();
float3 P = optixGetObjectRayOrigin();
@@ -288,6 +286,24 @@ extern "C" __global__ void __intersection__curve()
__float_as_int(isect.u), // Attribute_0
__float_as_int(isect.v)); // Attribute_1
}
+
+}
+
+extern "C" __global__ void __intersection__curve_ribbon()
+{
+ const uint prim = optixGetPrimitiveIndex();
+ const uint type = kernel_tex_fetch(__prim_type, prim);
+
+ if (type & (PRIMITIVE_CURVE_RIBBON | PRIMITIVE_MOTION_CURVE_RIBBON)) {
+ optix_intersection_curve(prim, type);
+ }
+}
+
+extern "C" __global__ void __intersection__curve_all()
+{
+ const uint prim = optixGetPrimitiveIndex();
+ const uint type = kernel_tex_fetch(__prim_type, prim);
+ optix_intersection_curve(prim, type);
}
#endif