[Bf-blender-cvs] [f613c4c0953] master: Cycles: MetalRT support (kernel side)
Michael Jones
noreply at git.blender.org
Mon Nov 29 16:20:33 CET 2021
Commit: f613c4c0953ebaf993ecd55b12bab9cf2196dac4
Author: Michael Jones
Date: Mon Nov 29 15:06:22 2021 +0000
Branches: master
https://developer.blender.org/rBf613c4c0953ebaf993ecd55b12bab9cf2196dac4
Cycles: MetalRT support (kernel side)
This patch adds MetalRT support to Cycles kernel code. It is mostly additive in nature or confined to Metal-specific code, however there are a few areas where this interacts with other code:
- MetalRT closely follows the Optix implementation, and in some cases (notably handling of transforms) it makes sense to extend Optix special-casing to MetalRT. For these generalisations we now have `__KERNEL_GPU_RAYTRACING__` instead of `__KERNEL_OPTIX__`.
- MetalRT doesn't support primitive offsetting (as with `primitiveIndexOffset` in Optix), so we define and populate a new kernel texture, `__object_prim_offset`, containing per-object primitive / curve-segment offsets. This is referenced and applied in MetalRT intersection handlers.
- Two new BVH layout enum values have been added: `BVH_LAYOUT_METAL` and `BVH_LAYOUT_MULTI_METAL_EMBREE` for XPU mode). Some host-side enum case handling has been updated where it is trivial to do so.
Ref T92212
Reviewed By: brecht
Maniphest Tasks: T92212
Differential Revision: https://developer.blender.org/D13353
===================================================================
M intern/cycles/bvh/bvh.cpp
M intern/cycles/device/cpu/device_impl.cpp
M intern/cycles/device/multi/device.cpp
M intern/cycles/kernel/CMakeLists.txt
M intern/cycles/kernel/bvh/bvh.h
A intern/cycles/kernel/bvh/metal.h
M intern/cycles/kernel/device/gpu/kernel.h
M intern/cycles/kernel/device/metal/compat.h
M intern/cycles/kernel/device/metal/context_begin.h
M intern/cycles/kernel/device/metal/kernel.metal
M intern/cycles/kernel/device/optix/compat.h
M intern/cycles/kernel/geom/motion_triangle_intersect.h
M intern/cycles/kernel/geom/triangle_intersect.h
M intern/cycles/kernel/integrator/subsurface_disk.h
M intern/cycles/kernel/integrator/subsurface_random_walk.h
M intern/cycles/kernel/textures.h
M intern/cycles/kernel/types.h
M intern/cycles/scene/geometry.cpp
M intern/cycles/scene/object.cpp
M intern/cycles/scene/object.h
M intern/cycles/scene/scene.cpp
M intern/cycles/scene/scene.h
M intern/cycles/util/math_float3.h
M intern/cycles/util/transform.h
===================================================================
diff --git a/intern/cycles/bvh/bvh.cpp b/intern/cycles/bvh/bvh.cpp
index ae6655eb27b..d3c8e4db6d0 100644
--- a/intern/cycles/bvh/bvh.cpp
+++ b/intern/cycles/bvh/bvh.cpp
@@ -40,8 +40,11 @@ const char *bvh_layout_name(BVHLayout layout)
return "EMBREE";
case BVH_LAYOUT_OPTIX:
return "OPTIX";
+ case BVH_LAYOUT_METAL:
+ return "METAL";
case BVH_LAYOUT_MULTI_OPTIX:
case BVH_LAYOUT_MULTI_OPTIX_EMBREE:
+ case BVH_LAYOUT_MULTI_METAL_EMBREE:
return "MULTI";
case BVH_LAYOUT_ALL:
return "ALL";
@@ -105,7 +108,10 @@ BVH *BVH::create(const BVHParams ¶ms,
#endif
case BVH_LAYOUT_MULTI_OPTIX:
case BVH_LAYOUT_MULTI_OPTIX_EMBREE:
+ case BVH_LAYOUT_MULTI_METAL_EMBREE:
return new BVHMulti(params, geometry, objects);
+ case BVH_LAYOUT_METAL:
+ /* host-side changes for BVH_LAYOUT_METAL are imminent */
case BVH_LAYOUT_NONE:
case BVH_LAYOUT_ALL:
break;
diff --git a/intern/cycles/device/cpu/device_impl.cpp b/intern/cycles/device/cpu/device_impl.cpp
index 2ad76de70ca..62b9cc93dae 100644
--- a/intern/cycles/device/cpu/device_impl.cpp
+++ b/intern/cycles/device/cpu/device_impl.cpp
@@ -274,7 +274,8 @@ void CPUDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
{
#ifdef WITH_EMBREE
if (bvh->params.bvh_layout == BVH_LAYOUT_EMBREE ||
- bvh->params.bvh_layout == BVH_LAYOUT_MULTI_OPTIX_EMBREE) {
+ bvh->params.bvh_layout == BVH_LAYOUT_MULTI_OPTIX_EMBREE ||
+ bvh->params.bvh_layout == BVH_LAYOUT_MULTI_METAL_EMBREE) {
BVHEmbree *const bvh_embree = static_cast<BVHEmbree *>(bvh);
if (refit) {
bvh_embree->refit(progress);
diff --git a/intern/cycles/device/multi/device.cpp b/intern/cycles/device/multi/device.cpp
index e319246d4f4..2513df63489 100644
--- a/intern/cycles/device/multi/device.cpp
+++ b/intern/cycles/device/multi/device.cpp
@@ -129,6 +129,10 @@ class MultiDevice : public Device {
if ((bvh_layout_mask_all & BVH_LAYOUT_OPTIX_EMBREE) == BVH_LAYOUT_OPTIX_EMBREE) {
return BVH_LAYOUT_MULTI_OPTIX_EMBREE;
}
+ const BVHLayoutMask BVH_LAYOUT_METAL_EMBREE = (BVH_LAYOUT_METAL | BVH_LAYOUT_EMBREE);
+ if ((bvh_layout_mask_all & BVH_LAYOUT_METAL_EMBREE) == BVH_LAYOUT_METAL_EMBREE) {
+ return BVH_LAYOUT_MULTI_METAL_EMBREE;
+ }
return bvh_layout_mask;
}
@@ -151,7 +155,8 @@ class MultiDevice : public Device {
}
assert(bvh->params.bvh_layout == BVH_LAYOUT_MULTI_OPTIX ||
- bvh->params.bvh_layout == BVH_LAYOUT_MULTI_OPTIX_EMBREE);
+ bvh->params.bvh_layout == BVH_LAYOUT_MULTI_OPTIX_EMBREE ||
+ bvh->params.bvh_layout == BVH_LAYOUT_MULTI_METAL_EMBREE);
BVHMulti *const bvh_multi = static_cast<BVHMulti *>(bvh);
bvh_multi->sub_bvhs.resize(devices.size());
diff --git a/intern/cycles/kernel/CMakeLists.txt b/intern/cycles/kernel/CMakeLists.txt
index d759399b04d..674eb702814 100644
--- a/intern/cycles/kernel/CMakeLists.txt
+++ b/intern/cycles/kernel/CMakeLists.txt
@@ -207,6 +207,7 @@ set(SRC_KERNEL_BVH_HEADERS
bvh/volume.h
bvh/volume_all.h
bvh/embree.h
+ bvh/metal.h
)
set(SRC_KERNEL_CAMERA_HEADERS
diff --git a/intern/cycles/kernel/bvh/bvh.h b/intern/cycles/kernel/bvh/bvh.h
index 0e083812355..33d2e44471a 100644
--- a/intern/cycles/kernel/bvh/bvh.h
+++ b/intern/cycles/kernel/bvh/bvh.h
@@ -31,6 +31,10 @@
# include "kernel/bvh/embree.h"
#endif
+#ifdef __METALRT__
+# include "kernel/bvh/metal.h"
+#endif
+
#include "kernel/bvh/types.h"
#include "kernel/bvh/util.h"
@@ -38,7 +42,7 @@
CCL_NAMESPACE_BEGIN
-#ifndef __KERNEL_OPTIX__
+#if !defined(__KERNEL_GPU_RAYTRACING__)
/* Regular BVH traversal */
@@ -139,7 +143,7 @@ CCL_NAMESPACE_BEGIN
# undef BVH_NAME_EVAL
# undef BVH_FUNCTION_FULL_NAME
-#endif /* __KERNEL_OPTIX__ */
+#endif /* !defined(__KERNEL_GPU_RAYTRACING__) */
ccl_device_inline bool scene_intersect_valid(ccl_private const Ray *ray)
{
@@ -205,7 +209,95 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg,
isect->type = p5;
return p5 != PRIMITIVE_NONE;
-#else /* __KERNEL_OPTIX__ */
+#elif defined(__METALRT__)
+
+ if (!scene_intersect_valid(ray)) {
+ isect->t = ray->t;
+ isect->type = PRIMITIVE_NONE;
+ return false;
+ }
+
+# if defined(__KERNEL_DEBUG__)
+ if (is_null_instance_acceleration_structure(metal_ancillaries->accel_struct)) {
+ isect->t = ray->t;
+ isect->type = PRIMITIVE_NONE;
+ kernel_assert(!"Invalid metal_ancillaries->accel_struct pointer");
+ return false;
+ }
+
+ if (is_null_intersection_function_table(metal_ancillaries->ift_default)) {
+ isect->t = ray->t;
+ isect->type = PRIMITIVE_NONE;
+ kernel_assert(!"Invalid ift_default");
+ return false;
+ }
+# endif
+
+ metal::raytracing::ray r(ray->P, ray->D, 0.0f, ray->t);
+ metalrt_intersector_type metalrt_intersect;
+
+ if (!kernel_data.bvh.have_curves) {
+ metalrt_intersect.assume_geometry_type(metal::raytracing::geometry_type::triangle);
+ }
+
+ MetalRTIntersectionPayload payload;
+ payload.u = 0.0f;
+ payload.v = 0.0f;
+ payload.visibility = visibility;
+
+ typename metalrt_intersector_type::result_type intersection;
+
+ uint ray_mask = visibility & 0xFF;
+ if (0 == ray_mask && (visibility & ~0xFF) != 0) {
+ ray_mask = 0xFF;
+ /* No further intersector setup required: Default MetalRT behaviour is anyhit */
+ }
+ else if (visibility & PATH_RAY_SHADOW_OPAQUE) {
+ /* No further intersector setup required: Shadow ray early termination is controlled by the
+ * intersection handler */
+ }
+
+# if defined(__METALRT_MOTION__)
+ payload.time = ray->time;
+ intersection = metalrt_intersect.intersect(r,
+ metal_ancillaries->accel_struct,
+ ray_mask,
+ ray->time,
+ metal_ancillaries->ift_default,
+ payload);
+# else
+ intersection = metalrt_intersect.intersect(
+ r, metal_ancillaries->accel_struct, ray_mask, metal_ancillaries->ift_default, payload);
+# endif
+
+ if (intersection.type == intersection_type::none) {
+ isect->t = ray->t;
+ isect->type = PRIMITIVE_NONE;
+
+ return false;
+ }
+
+ isect->t = intersection.distance;
+
+ isect->prim = payload.prim;
+ isect->type = payload.type;
+ isect->object = intersection.user_instance_id;
+
+ isect->t = intersection.distance;
+ if (intersection.type == intersection_type::triangle) {
+ isect->u = 1.0f - intersection.triangle_barycentric_coord.y -
+ intersection.triangle_barycentric_coord.x;
+ isect->v = intersection.triangle_barycentric_coord.x;
+ }
+ else {
+ isect->u = payload.u;
+ isect->v = payload.v;
+ }
+
+ return isect->type != PRIMITIVE_NONE;
+
+#else
+
if (!scene_intersect_valid(ray)) {
return false;
}
@@ -289,7 +381,69 @@ ccl_device_intersect bool scene_intersect_local(KernelGlobals kg,
p5);
return p5;
-# else /* __KERNEL_OPTIX__ */
+# elif defined(__METALRT__)
+ if (!scene_intersect_valid(ray)) {
+ if (local_isect) {
+ local_isect->num_hits = 0;
+ }
+ return false;
+ }
+
+# if defined(__KERNEL_DEBUG__)
+ if (is_null_instance_acceleration_structure(metal_ancillaries->accel_struct)) {
+ if (local_isect) {
+ local_isect->num_hits = 0;
+ }
+ kernel_assert(!"Invalid metal_ancillaries->accel_struct pointer");
+ return false;
+ }
+
+ if (is_null_intersection_function_table(metal_ancillaries->ift_local)) {
+ if (local_isect) {
+ local_isect->num_hits = 0;
+ }
+ kernel_assert(!"Invalid ift_local");
+ return false;
+ }
+# endif
+
+ metal::raytracing::ray r(ray->P, ray->D, 0.0f, ray->t);
+ metalrt_intersector_type metalrt_intersect;
+
+ metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque);
+ if (!kernel_data.bvh.have_curves) {
+ metalrt_intersect.assume_geometry_type(metal::raytracing::geometry_type::triangle);
+ }
+
+ MetalRTIntersectionLocalPayload payload;
+ payload.local_object = local_object;
+ payload.max_hits = max_hits;
+ payload.local_isect.num_hits = 0;
+ if (lcg_state) {
+ payload.has_lcg_state = true;
+ payload.lcg_state = *lcg_state;
+ }
+ payload.result = false;
+
+ typename metalrt_intersector_type::result_type intersection;
+
+# if defined(__METALRT_MOTION__)
+ intersection = metalrt_intersect.intersect(
+ r, metal_ancillaries->accel_struct, 0xFF, ray->time, metal_ancillaries->ift_local, payload);
+# else
+ intersection = metalrt_intersect.intersect(
+ r, metal_ancillaries->accel_struct, 0xFF, metal_ancillaries->ift_local, payload);
+# endif
+
+ if (lcg_state) {
+ *lcg_state = payload.lcg_state;
+ }
+ *local_isect = payload.local_isect;
+
+ return payload.result;
+
+# else
+
if (!scene_intersect_valid(ray)) {
if (local_isect) {
local_isect->num_hits = 0;
@@ -406,7 +560,67 @@ ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg,
*throughput = __uint_as_float(p1);
return p5;
-# else /* __KERNEL_OPTIX__ */
+# elif defined(__METALRT__)
+
+ if (!scene_intersect_valid(ray)) {
+ return false;
+ }
+
+# if defined(__KERNEL_DEBUG__)
+ if (is_null_instance_acceleration_structure(metal_ancillaries->accel_struct)) {
+ kernel_assert(!"Invalid metal_ancillaries->accel_struct pointer");
+ return false;
+ }
+
+ if (is_null_intersection_function_table(metal_ancillaries->ift_shadow)) {
+ kernel_assert(!"Invalid ift_shadow");
+ return false;
+ }
+# endif
+
+ metal::raytracing::ray r(ray->P, ray->D, 0.0f, ray->t);
+ metalrt_intersector_type metalrt_intersect;
+
+ metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque);
+ if (!kernel_data.bvh.have_curves) {
+ metalrt_intersect.assume_geometry_type(metal::raytracing::geometry_type::triangle);
+ }
+
+ MetalRTIntersectionShadowPayload payload;
+ payload.visibility = visibility;
+ payload.max_hits = max_hits;
+ payload.num_hits = 0;
+ payload.num_recorded_hits = 0;
+ payload.throughput = 1.0f;
+ payload.result = false;
+ payload.state = state;
+
+ uint ray_
@@ Diff output truncated at 10240 characters. @@
More information about the Bf-blender-cvs
mailing list