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:
authorPatrick Mours <pmours@nvidia.com>2020-07-06 13:25:54 +0300
committerPatrick Mours <pmours@nvidia.com>2020-07-07 16:39:02 +0300
commit737bd549b6eeee81f0573ad8e305fb8d888d82ec (patch)
treefc8aaca871557d652287cb3d4f9ee3033ff06e46 /intern/cycles/kernel
parent95f0f312799e10e4a9f5c884f8f20ec76d0ff363 (diff)
Cycles: Add support for native OptiX curve primitive
This patch adds support for the curve primitive from OptiX to Cycles. It's currently hidden behind a debug option, since there can be some slight rendering differences still (because no backface culling is performed and something seems off with endcaps). The curve primitive was added with the OptiX 7.1 SDK and requires a r450 driver or newer, so this also updates the codebase to be able to build with the new SDK. Reviewed By: brecht Differential Revision: https://developer.blender.org/D8223
Diffstat (limited to 'intern/cycles/kernel')
-rw-r--r--intern/cycles/kernel/bvh/bvh.h23
-rw-r--r--intern/cycles/kernel/geom/geom_curve_intersect.h9
-rw-r--r--intern/cycles/kernel/kernels/optix/kernel_optix.cu48
3 files changed, 49 insertions, 31 deletions
diff --git a/intern/cycles/kernel/bvh/bvh.h b/intern/cycles/kernel/bvh/bvh.h
index 80b58f46329..3049f243ae9 100644
--- a/intern/cycles/kernel/bvh/bvh.h
+++ b/intern/cycles/kernel/bvh/bvh.h
@@ -172,11 +172,11 @@ ccl_device_intersect bool scene_intersect(KernelGlobals *kg,
0.0f,
ray->t,
ray->time,
- 0xFF,
+ 0xF,
OPTIX_RAY_FLAG_NONE,
+ 0, // SBT offset for PG_HITD
0,
0,
- 0, // SBT offset for PG_HITD
p0,
p1,
p2,
@@ -264,12 +264,13 @@ ccl_device_intersect bool scene_intersect_local(KernelGlobals *kg,
0.0f,
ray->t,
ray->time,
+ // Skip curves
+ 0x3,
// Need to always call into __anyhit__kernel_optix_local_hit
- 0xFF,
OPTIX_RAY_FLAG_ENFORCE_ANYHIT,
- 1,
+ 2, // SBT offset for PG_HITL
+ 0,
0,
- 0, // SBT offset for PG_HITL
p0,
p1,
p2,
@@ -374,12 +375,12 @@ ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals *kg,
0.0f,
ray->t,
ray->time,
+ 0xF,
// Need to always call into __anyhit__kernel_optix_shadow_all_hit
- 0xFF,
OPTIX_RAY_FLAG_ENFORCE_ANYHIT,
- 2,
+ 1, // SBT offset for PG_HITS
+ 0,
0,
- 0, // SBT offset for PG_HITS
p0,
p1,
*num_hits,
@@ -458,12 +459,12 @@ ccl_device_intersect bool scene_intersect_volume(KernelGlobals *kg,
0.0f,
ray->t,
ray->time,
- // Visibility mask set to only intersect objects with volumes
- 0x02,
+ // Skip everything but volumes
+ 0x2,
OPTIX_RAY_FLAG_NONE,
+ 0, // SBT offset for PG_HITD
0,
0,
- 0, // SBT offset for PG_HITD
p0,
p1,
p2,
diff --git a/intern/cycles/kernel/geom/geom_curve_intersect.h b/intern/cycles/kernel/geom/geom_curve_intersect.h
index c04dbee52cc..06d2c016f5b 100644
--- a/intern/cycles/kernel/geom/geom_curve_intersect.h
+++ b/intern/cycles/kernel/geom/geom_curve_intersect.h
@@ -734,7 +734,6 @@ ccl_device_inline void curve_shader_setup(KernelGlobals *kg,
}
sd->u = isect->u;
- sd->v = isect->v;
P = P + D * t;
@@ -750,6 +749,7 @@ ccl_device_inline void curve_shader_setup(KernelGlobals *kg,
sd->N = normalize(sine * bitangent - cosine * normalize(cross(tangent, bitangent)));
sd->Ng = -D;
+ sd->v = isect->v;
# if 0
/* This approximates the position and geometric normal of a thick curve too,
@@ -764,8 +764,11 @@ ccl_device_inline void curve_shader_setup(KernelGlobals *kg,
* This could be optimized by recording the normal in the intersection,
* however for Optix this would go beyond the size of the payload. */
const float3 P_inside = float4_to_float3(catmull_rom_basis_eval(P_curve, isect->u));
- sd->Ng = normalize(P - P_inside);
- sd->N = sd->Ng;
+ const float3 Ng = normalize(P - P_inside);
+
+ sd->N = Ng;
+ sd->Ng = Ng;
+ sd->v = 0.0f;
}
# ifdef __DPDU__
diff --git a/intern/cycles/kernel/kernels/optix/kernel_optix.cu b/intern/cycles/kernel/kernels/optix/kernel_optix.cu
index c730d952ed4..3b166e59dfd 100644
--- a/intern/cycles/kernel/kernels/optix/kernel_optix.cu
+++ b/intern/cycles/kernel/kernels/optix/kernel_optix.cu
@@ -15,6 +15,7 @@
* limitations under the License.
*/
+// clang-format off
#include "kernel/kernel_compat_optix.h"
#include "util/util_atomic.h"
#include "kernel/kernel_types.h"
@@ -23,6 +24,7 @@
#include "kernel/kernel_path.h"
#include "kernel/kernel_bake.h"
+// clang-format on
template<typename T> ccl_device_forceinline T *get_payload_ptr_0()
{
@@ -139,8 +141,8 @@ extern "C" __global__ void __anyhit__kernel_optix_local_hit()
}
else {
if (local_isect->num_hits && optixGetRayTmax() > local_isect->hits[0].t) {
- // Record closest intersection only (do not terminate ray here, since there is no guarantee
- // about distance ordering in anyhit)
+ // Record closest intersection only
+ // Do not terminate ray here, since there is no guarantee about distance ordering in any-hit
return optixIgnoreIntersection();
}
@@ -153,15 +155,9 @@ extern "C" __global__ void __anyhit__kernel_optix_local_hit()
isect->object = get_object_id();
isect->type = kernel_tex_fetch(__prim_type, isect->prim);
- if (optixIsTriangleHit()) {
- const float2 barycentrics = optixGetTriangleBarycentrics();
- isect->u = 1.0f - barycentrics.y - barycentrics.x;
- isect->v = barycentrics.x;
- }
- else {
- isect->u = __uint_as_float(optixGetAttribute_0());
- isect->v = __uint_as_float(optixGetAttribute_1());
- }
+ const float2 barycentrics = optixGetTriangleBarycentrics();
+ isect->u = 1.0f - barycentrics.y - barycentrics.x;
+ isect->v = barycentrics.x;
// Record geometric normal
const uint tri_vindex = kernel_tex_fetch(__prim_tri_index, isect->prim);
@@ -198,10 +194,18 @@ extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit()
isect->u = 1.0f - barycentrics.y - barycentrics.x;
isect->v = barycentrics.x;
}
+# ifdef __HAIR__
else {
- isect->u = __uint_as_float(optixGetAttribute_0());
+ const float u = __uint_as_float(optixGetAttribute_0());
+ isect->u = u;
isect->v = __uint_as_float(optixGetAttribute_1());
+
+ // Filter out curve endcaps
+ if (u == 0.0f || u == 1.0f) {
+ return optixIgnoreIntersection();
+ }
}
+# endif
# ifdef __TRANSPARENT_SHADOWS__
// Detect if this surface has a shader with transparent shadows
@@ -213,7 +217,6 @@ extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit()
# ifdef __TRANSPARENT_SHADOWS__
}
- // TODO(pmours): Do we need REQUIRE_UNIQUE_ANYHIT for this to work?
optixSetPayload_2(optixGetPayload_2() + 1); // num_hits++
// Continue tracing
@@ -227,13 +230,25 @@ extern "C" __global__ void __anyhit__kernel_optix_visibility_test()
uint visibility = optixGetPayload_4();
#ifdef __VISIBILITY_FLAG__
const uint prim = optixGetPrimitiveIndex();
- if ((kernel_tex_fetch(__prim_visibility, prim) & visibility) == 0)
+ if ((kernel_tex_fetch(__prim_visibility, prim) & visibility) == 0) {
return optixIgnoreIntersection();
+ }
+#endif
+
+#ifdef __HAIR__
+ if (!optixIsTriangleHit()) {
+ // Filter out curve endcaps
+ const float u = __uint_as_float(optixGetAttribute_0());
+ if (u == 0.0f || u == 1.0f) {
+ return optixIgnoreIntersection();
+ }
+ }
#endif
// Shadow ray early termination
- if (visibility & PATH_RAY_SHADOW_OPAQUE)
+ if (visibility & PATH_RAY_SHADOW_OPAQUE) {
return optixTerminateRay();
+ }
}
extern "C" __global__ void __closesthit__kernel_optix_hit()
@@ -250,7 +265,7 @@ extern "C" __global__ void __closesthit__kernel_optix_hit()
optixSetPayload_2(__float_as_uint(barycentrics.x));
}
else {
- optixSetPayload_1(optixGetAttribute_0());
+ optixSetPayload_1(optixGetAttribute_0()); // Same as 'optixGetCurveParameter()'
optixSetPayload_2(optixGetAttribute_1());
}
}
@@ -286,7 +301,6 @@ ccl_device_inline void optix_intersection_curve(const uint prim, const uint type
__float_as_int(isect.u), // Attribute_0
__float_as_int(isect.v)); // Attribute_1
}
-
}
extern "C" __global__ void __intersection__curve_ribbon()