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

Cycles: Use new MetalRT curve primitives for 3D curves and ribbons

Browse Source 
This patch updates the experimental MetalRT code path to use new [curve primitives](https://developer.apple.com/videos/play/wwdc2023/10128/) which were recently added in macOS 14. This replaces the previous custom box intersection implementation, allowing the driver to better optimise curve acceleration structures for the GPU. On existing hardware, this can speed up MetalRT renders by up to 40% for scenes that use hair / curve primitives extensively.

The MetalRT option will only be available on macOS >= 14, and requires Xcode >= 15 to build (otherwise the option will be compiled out).

Authored by Marco Giordano, Michael Jones, and Jason Fielder

---
Before / after render times (M1 Max MacBook Pro, macOS 14 beta, MetalRT enabled):
```
                  Custom box intersection      MetalRT curve primitives       Speedup
fishy_cat           111.5                         80.5                         1.39
koro                114.4                         86.7                         1.32
sinosauropteryx     291.8                        279.2                         1.05
spring              142.3                        142.2                         1.00
victor              442.7                        347.7                         1.27
```

---

Pull Request: https://projects.blender.org/blender/blender/pulls/111795
This commit is contained in:
Michael Jones
2023-09-13 16:02:49 +02:00
committed by Michael Jones (Apple)
parent 36f31f1eff
commit 6c98cb73ac
13 changed files with 706 additions and 568 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -1681,9 +1681,13 @@ class CyclesPreferences(bpy.types.AddonPreferences):
import _cycles
has_peer_memory = 0
has_rt_api_support = False
for device in _cycles.available_devices(compute_device_type):
if device[3] and self.find_existing_device_entry(device).use:
has_peer_memory += 1
if device[4] and self.find_existing_device_entry(device).use:
has_rt_api_support = True
if has_peer_memory > 1:
row = layout.row()
row.use_property_split = True
@@ -1700,13 +1704,14 @@ class CyclesPreferences(bpy.types.AddonPreferences):
# MetalRT only works on Apple Silicon and Navi2.
is_arm64 = platform.machine() == 'arm64'
if is_arm64 or is_navi_2:
if is_arm64 or (is_navi_2 and has_rt_api_support):
col = layout.column()
col.use_property_split = True
# Kernel specialization is only supported on Apple Silicon
if is_arm64:
col.prop(self, "kernel_optimization_level")
col.prop(self, "use_metalrt")
if has_rt_api_support:
col.prop(self, "use_metalrt")
if compute_device_type == 'HIP':
has_cuda, has_optix, has_hip, has_metal, has_oneapi, has_hiprt = _cycles.get_device_types()

3
intern/cycles/blender/python.cpp
View File

@@ -410,11 +410,12 @@ static PyObject *available_devices_func(PyObject * /*self*/, PyObject *args)
for (size_t i = 0; i < devices.size(); i++) {
DeviceInfo &device = devices[i];
string type_name = Device::string_from_type(device.type);
PyObject *device_tuple = PyTuple_New(4);
PyObject *device_tuple = PyTuple_New(5);
PyTuple_SET_ITEM(device_tuple, 0, pyunicode_from_string(device.description.c_str()));
PyTuple_SET_ITEM(device_tuple, 1, pyunicode_from_string(type_name.c_str()));
PyTuple_SET_ITEM(device_tuple, 2, pyunicode_from_string(device.id.c_str()));
PyTuple_SET_ITEM(device_tuple, 3, PyBool_FromLong(device.has_peer_memory));
PyTuple_SET_ITEM(device_tuple, 4, PyBool_FromLong(device.use_hardware_raytracing));
PyTuple_SET_ITEM(ret, i, device_tuple);
}

4
intern/cycles/device/metal/bvh.h
View File

@@ -22,7 +22,9 @@ class BVHMetal : public BVH {
API_AVAILABLE(macos(11.0))
vector<id<MTLAccelerationStructure>> blas_array;
vector<uint32_t> blas_lookup;
API_AVAILABLE(macos(11.0))
vector<id<MTLAccelerationStructure>> unique_blas_array;
bool motion_blur = false;

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

@@ -132,6 +132,7 @@ bool BVHMetal::build_BLAS_mesh(Progress &progress,
geomDescMotion.indexType = MTLIndexTypeUInt32;
geomDescMotion.triangleCount = num_indices / 3;
geomDescMotion.intersectionFunctionTableOffset = 0;
geomDescMotion.opaque = true;
geomDesc = geomDescMotion;
}
@@ -146,6 +147,7 @@ bool BVHMetal::build_BLAS_mesh(Progress &progress,
geomDescNoMotion.indexType = MTLIndexTypeUInt32;
geomDescNoMotion.triangleCount = num_indices / 3;
geomDescNoMotion.intersectionFunctionTableOffset = 0;
geomDescNoMotion.opaque = true;
geomDesc = geomDescNoMotion;
}
@@ -165,6 +167,7 @@ bool BVHMetal::build_BLAS_mesh(Progress &progress,
accelDesc.motionEndBorderMode = MTLMotionBorderModeClamp;
accelDesc.motionKeyframeCount = num_motion_steps;
}
accelDesc.usage |= MTLAccelerationStructureUsageExtendedLimits;
if (!use_fast_trace_bvh) {
accelDesc.usage |= (MTLAccelerationStructureUsageRefit |
@@ -255,7 +258,8 @@ bool BVHMetal::build_BLAS_hair(Progress &progress,
Geometry *const geom,
bool refit)
{
if (@available(macos 12.0, *)) {
# if defined(MAC_OS_VERSION_14_0)
if (@available(macos 14.0, *)) {
/* Build BLAS for hair curves */
Hair *hair = static_cast<Hair *>(geom);
if (hair->num_curves() == 0) {
@@ -268,7 +272,6 @@ bool BVHMetal::build_BLAS_hair(Progress &progress,
/*------------------------------------------------*/
const bool use_fast_trace_bvh = (params.bvh_type == BVH_TYPE_STATIC);
const size_t num_segments = hair->num_segments();
size_t num_motion_steps = 1;
Attribute *motion_keys = hair->attributes.find(ATTR_STD_MOTION_VERTEX_POSITION);
@@ -276,8 +279,6 @@ bool BVHMetal::build_BLAS_hair(Progress &progress,
num_motion_steps = hair->get_motion_steps();
}
const size_t num_aabbs = num_segments * num_motion_steps;
MTLResourceOptions storage_mode;
if (device.hasUnifiedMemory) {
storage_mode = MTLResourceStorageModeShared;
@@ -286,91 +287,197 @@ bool BVHMetal::build_BLAS_hair(Progress &progress,
storage_mode = MTLResourceStorageModeManaged;
}
/* Allocate a GPU buffer for the AABB data and populate it */
id<MTLBuffer> aabbBuf = [device
newBufferWithLength:num_aabbs * sizeof(MTLAxisAlignedBoundingBox)
options:storage_mode];
MTLAxisAlignedBoundingBox *aabb_data = (MTLAxisAlignedBoundingBox *)[aabbBuf contents];
/* Get AABBs for each motion step */
size_t center_step = (num_motion_steps - 1) / 2;
for (size_t step = 0; step < num_motion_steps; ++step) {
/* The center step for motion vertices is not stored in the attribute */
const float3 *keys = hair->get_curve_keys().data();
if (step != center_step) {
size_t attr_offset = (step > center_step) ? step - 1 : step;
/* Technically this is a float4 array, but sizeof(float3) == sizeof(float4) */
keys = motion_keys->data_float3() + attr_offset * hair->get_curve_keys().size();
}
for (size_t j = 0, i = 0; j < hair->num_curves(); ++j) {
const Hair::Curve curve = hair->get_curve(j);
for (int segment = 0; segment < curve.num_segments(); ++segment, ++i) {
{
BoundBox bounds = BoundBox::empty;
curve.bounds_grow(segment, keys, hair->get_curve_radius().data(), bounds);
const size_t index = step * num_segments + i;
aabb_data[index].min = (MTLPackedFloat3 &)bounds.min;
aabb_data[index].max = (MTLPackedFloat3 &)bounds.max;
}
}
}
}
if (storage_mode == MTLResourceStorageModeManaged) {
[aabbBuf didModifyRange:NSMakeRange(0, aabbBuf.length)];
}
# if 0
for (size_t i=0; i<num_aabbs && i < 400; i++) {
MTLAxisAlignedBoundingBox& bb = aabb_data[i];
printf(" %d: %.1f,%.1f,%.1f -- %.1f,%.1f,%.1f\n", int(i), bb.min.x, bb.min.y, bb.min.z, bb.max.x, bb.max.y, bb.max.z);
}
# endif
id<MTLBuffer> cpBuffer = nil;
id<MTLBuffer> radiusBuffer = nil;
id<MTLBuffer> idxBuffer = nil;
MTLAccelerationStructureGeometryDescriptor *geomDesc;
if (motion_blur) {
std::vector<MTLMotionKeyframeData *> aabb_ptrs;
aabb_ptrs.reserve(num_motion_steps);
MTLAccelerationStructureMotionCurveGeometryDescriptor *geomDescCrv =
[MTLAccelerationStructureMotionCurveGeometryDescriptor descriptor];
uint64_t numKeys = hair->num_keys();
uint64_t numCurves = hair->num_curves();
const array<float> &radiuses = hair->get_curve_radius();
/* Gather the curve geometry. */
std::vector<float3> cpData;
std::vector<int> idxData;
std::vector<float> radiusData;
cpData.reserve(numKeys);
radiusData.reserve(numKeys);
std::vector<int> step_offsets;
for (size_t step = 0; step < num_motion_steps; ++step) {
MTLMotionKeyframeData *k = [MTLMotionKeyframeData data];
k.buffer = aabbBuf;
k.offset = step * num_segments * sizeof(MTLAxisAlignedBoundingBox);
aabb_ptrs.push_back(k);
/* The center step for motion vertices is not stored in the attribute. */
const float3 *keys = hair->get_curve_keys().data();
size_t center_step = (num_motion_steps - 1) / 2;
if (step != center_step) {
size_t attr_offset = (step > center_step) ? step - 1 : step;
/* Technically this is a float4 array, but sizeof(float3) == sizeof(float4). */
keys = motion_keys->data_float3() + attr_offset * numKeys;
}
step_offsets.push_back(cpData.size());
for (int c = 0; c < numCurves; ++c) {
const Hair::Curve curve = hair->get_curve(c);
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]);
for (int s = 0; s < segCount; ++s) {
if (step == 0) {
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]);
}
}
MTLAccelerationStructureMotionBoundingBoxGeometryDescriptor *geomDescMotion =
[MTLAccelerationStructureMotionBoundingBoxGeometryDescriptor descriptor];
geomDescMotion.boundingBoxBuffers = [NSArray arrayWithObjects:aabb_ptrs.data()
count:aabb_ptrs.size()];
geomDescMotion.boundingBoxCount = num_segments;
geomDescMotion.boundingBoxStride = sizeof(aabb_data[0]);
geomDescMotion.intersectionFunctionTableOffset = 1;
/* Allocate and populate MTLBuffers for geometry. */
idxBuffer = [device newBufferWithBytes:idxData.data()
length:idxData.size() * sizeof(int)
options:storage_mode];
cpBuffer = [device newBufferWithBytes:cpData.data()
length:cpData.size() * sizeof(float3)
options:storage_mode];
radiusBuffer = [device newBufferWithBytes:radiusData.data()
length:radiusData.size() * sizeof(float)
options:storage_mode];
std::vector<MTLMotionKeyframeData *> cp_ptrs;
std::vector<MTLMotionKeyframeData *> radius_ptrs;
cp_ptrs.reserve(num_motion_steps);
radius_ptrs.reserve(num_motion_steps);
for (size_t step = 0; step < num_motion_steps; ++step) {
MTLMotionKeyframeData *k = [MTLMotionKeyframeData data];
k.buffer = cpBuffer;
k.offset = step_offsets[step] * sizeof(float3);
cp_ptrs.push_back(k);
k = [MTLMotionKeyframeData data];
k.buffer = radiusBuffer;
k.offset = step_offsets[step] * sizeof(float);
radius_ptrs.push_back(k);
}
if (storage_mode == MTLResourceStorageModeManaged) {
[cpBuffer didModifyRange:NSMakeRange(0, cpBuffer.length)];
[idxBuffer didModifyRange:NSMakeRange(0, idxBuffer.length)];
[radiusBuffer didModifyRange:NSMakeRange(0, radiusBuffer.length)];
}
geomDescCrv.controlPointBuffers = [NSArray arrayWithObjects:cp_ptrs.data()
count:cp_ptrs.size()];
geomDescCrv.radiusBuffers = [NSArray arrayWithObjects:radius_ptrs.data()
count:radius_ptrs.size()];
geomDescCrv.controlPointCount = cpData.size();
geomDescCrv.controlPointStride = sizeof(float3);
geomDescCrv.controlPointFormat = MTLAttributeFormatFloat3;
geomDescCrv.radiusStride = sizeof(float);
geomDescCrv.radiusFormat = MTLAttributeFormatFloat;
geomDescCrv.segmentCount = idxData.size();
geomDescCrv.segmentControlPointCount = 4;
geomDescCrv.curveType = (hair->curve_shape == CURVE_RIBBON) ? MTLCurveTypeFlat :
MTLCurveTypeRound;
geomDescCrv.curveBasis = MTLCurveBasisCatmullRom;
geomDescCrv.curveEndCaps = MTLCurveEndCapsDisk;
geomDescCrv.indexType = MTLIndexTypeUInt32;
geomDescCrv.indexBuffer = idxBuffer;
geomDescCrv.intersectionFunctionTableOffset = 1;
/* Force a single any-hit call, so shadow record-all behavior works correctly */
/* (Match optix behavior: unsigned int build_flags =
* OPTIX_GEOMETRY_FLAG_REQUIRE_SINGLE_ANYHIT_CALL;) */
geomDescMotion.allowDuplicateIntersectionFunctionInvocation = false;
geomDescMotion.opaque = true;
geomDesc = geomDescMotion;
geomDescCrv.allowDuplicateIntersectionFunctionInvocation = false;
geomDescCrv.opaque = true;
geomDesc = geomDescCrv;
}
else {
MTLAccelerationStructureBoundingBoxGeometryDescriptor *geomDescNoMotion =
[MTLAccelerationStructureBoundingBoxGeometryDescriptor descriptor];
geomDescNoMotion.boundingBoxBuffer = aabbBuf;
geomDescNoMotion.boundingBoxBufferOffset = 0;
geomDescNoMotion.boundingBoxCount = int(num_aabbs);
geomDescNoMotion.boundingBoxStride = sizeof(aabb_data[0]);
geomDescNoMotion.intersectionFunctionTableOffset = 1;
MTLAccelerationStructureCurveGeometryDescriptor *geomDescCrv =
[MTLAccelerationStructureCurveGeometryDescriptor descriptor];
uint64_t numKeys = hair->num_keys();
uint64_t numCurves = hair->num_curves();
const array<float> &radiuses = hair->get_curve_radius();
/* Gather the curve geometry. */
std::vector<float3> cpData;
std::vector<int> idxData;
std::vector<float> radiusData;
cpData.reserve(numKeys);
radiusData.reserve(numKeys);
auto keys = hair->get_curve_keys();
for (int c = 0; c < numCurves; ++c) {
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]);
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]);
}
/* Allocate and populate MTLBuffers for geometry. */
idxBuffer = [device newBufferWithBytes:idxData.data()
length:idxData.size() * sizeof(int)
options:storage_mode];
cpBuffer = [device newBufferWithBytes:cpData.data()
length:cpData.size() * sizeof(float3)
options:storage_mode];
radiusBuffer = [device newBufferWithBytes:radiusData.data()
length:radiusData.size() * sizeof(float)
options:storage_mode];
if (storage_mode == MTLResourceStorageModeManaged) {
[cpBuffer didModifyRange:NSMakeRange(0, cpBuffer.length)];
[idxBuffer didModifyRange:NSMakeRange(0, idxBuffer.length)];
[radiusBuffer didModifyRange:NSMakeRange(0, radiusBuffer.length)];
}
geomDescCrv.controlPointBuffer = cpBuffer;
geomDescCrv.radiusBuffer = radiusBuffer;
geomDescCrv.controlPointCount = cpData.size();
geomDescCrv.controlPointStride = sizeof(float3);
geomDescCrv.controlPointFormat = MTLAttributeFormatFloat3;
geomDescCrv.controlPointBufferOffset = 0;
geomDescCrv.segmentCount = idxData.size();
geomDescCrv.segmentControlPointCount = 4;
geomDescCrv.curveType = (hair->curve_shape == CURVE_RIBBON) ? MTLCurveTypeFlat :
MTLCurveTypeRound;
geomDescCrv.curveBasis = MTLCurveBasisCatmullRom;
geomDescCrv.curveEndCaps = MTLCurveEndCapsDisk;
geomDescCrv.indexType = MTLIndexTypeUInt32;
geomDescCrv.indexBuffer = idxBuffer;
geomDescCrv.intersectionFunctionTableOffset = 1;
/* Force a single any-hit call, so shadow record-all behavior works correctly */
/* (Match optix behavior: unsigned int build_flags =
* OPTIX_GEOMETRY_FLAG_REQUIRE_SINGLE_ANYHIT_CALL;) */
geomDescNoMotion.allowDuplicateIntersectionFunctionInvocation = false;
geomDescNoMotion.opaque = true;
geomDesc = geomDescNoMotion;
geomDescCrv.allowDuplicateIntersectionFunctionInvocation = false;
geomDescCrv.opaque = true;
geomDesc = geomDescCrv;
}
MTLPrimitiveAccelerationStructureDescriptor *accelDesc =
@@ -389,6 +496,7 @@ bool BVHMetal::build_BLAS_hair(Progress &progress,
accelDesc.usage |= (MTLAccelerationStructureUsageRefit |
MTLAccelerationStructureUsagePreferFastBuild);
}
accelDesc.usage |= MTLAccelerationStructureUsageExtendedLimits;
MTLAccelerationStructureSizes accelSizes = [device
accelerationStructureSizesWithDescriptor:accelDesc];
@@ -423,10 +531,11 @@ bool BVHMetal::build_BLAS_hair(Progress &progress,
[accelCommands addCompletedHandler:^(id<MTLCommandBuffer> /*command_buffer*/) {
/* free temp resources */
[scratchBuf release];
[aabbBuf release];
[cpBuffer release];
[radiusBuffer release];
[idxBuffer release];
if (use_fast_trace_bvh) {
/* Compact the accel structure */
uint64_t compressed_size = *(uint64_t *)sizeBuf.contents;
dispatch_async(dispatch_get_global_queue(DISPATCH_QUEUE_PRIORITY_DEFAULT, 0), ^{
@@ -461,8 +570,10 @@ bool BVHMetal::build_BLAS_hair(Progress &progress,
accel_struct_building = true;
[accelCommands commit];
return true;
}
# endif /* MAC_OS_VERSION_14_0 */
return false;
}
@@ -605,10 +716,11 @@ bool BVHMetal::build_BLAS_pointcloud(Progress &progress,
if (motion_blur) {
accelDesc.motionStartTime = 0.0f;
accelDesc.motionEndTime = 1.0f;
accelDesc.motionStartBorderMode = MTLMotionBorderModeVanish;
accelDesc.motionEndBorderMode = MTLMotionBorderModeVanish;
// accelDesc.motionStartBorderMode = MTLMotionBorderModeVanish;
// accelDesc.motionEndBorderMode = MTLMotionBorderModeVanish;
accelDesc.motionKeyframeCount = num_motion_steps;
}
accelDesc.usage |= MTLAccelerationStructureUsageExtendedLimits;
if (!use_fast_trace_bvh) {
accelDesc.usage |= (MTLAccelerationStructureUsageRefit |
@@ -756,10 +868,11 @@ bool BVHMetal::build_TLAS(Progress &progress,
uint32_t num_instances = 0;
uint32_t num_motion_transforms = 0;
for (Object *ob : objects) {
/* Skip non-traceable objects */
num_instances++;
/* Skip motion for non-traceable objects */
if (!ob->is_traceable())
continue;
num_instances++;
if (ob->use_motion()) {
num_motion_transforms += max((size_t)1, ob->get_motion().size());
@@ -829,28 +942,40 @@ bool BVHMetal::build_TLAS(Progress &progress,
uint32_t instance_index = 0;
uint32_t motion_transform_index = 0;
// allocate look up buffer for wost case scenario
uint64_t count = objects.size();
blas_lookup.resize(count);
blas_array.clear();
blas_array.reserve(num_instances);
for (Object *ob : objects) {
/* Skip non-traceable objects */
if (!ob->is_traceable())
continue;
Geometry const *geom = ob->get_geometry();
BVHMetal const *blas = static_cast<BVHMetal const *>(geom->bvh);
if (!blas || !blas->accel_struct) {
/* Place a degenerate instance, to ensure [[instance_id]] equals ob->get_device_index()
* in our intersection functions */
if (motion_blur) {
MTLAccelerationStructureMotionInstanceDescriptor *instances =
(MTLAccelerationStructureMotionInstanceDescriptor *)[instanceBuf contents];
MTLAccelerationStructureMotionInstanceDescriptor &desc = instances[instance_index++];
memset(&desc, 0x00, sizeof(desc));
}
else {
MTLAccelerationStructureUserIDInstanceDescriptor *instances =
(MTLAccelerationStructureUserIDInstanceDescriptor *)[instanceBuf contents];
MTLAccelerationStructureUserIDInstanceDescriptor &desc = instances[instance_index++];
memset(&desc, 0x00, sizeof(desc));
}
blas_array.push_back(nil);
continue;
}
blas_array.push_back(blas->accel_struct);
uint32_t accel_struct_index = get_blas_index(blas);
/* Add some of the object visibility bits to the mask.
* __prim_visibility contains the combined visibility bits of all instances, so is not
* reliable if they differ between instances.
*
* METAL_WIP: OptiX visibility mask can only contain 8 bits, so have to trade-off here
* and select just a few important ones.
*/
uint32_t mask = ob->visibility_for_tracing() & 0xFF;
uint32_t mask = ob->visibility_for_tracing();
/* Have to have at least one bit in the mask, or else instance would always be culled. */
if (0 == mask) {
@@ -858,11 +983,25 @@ bool BVHMetal::build_TLAS(Progress &progress,
}
/* Set user instance ID to object index */
int object_index = ob->get_device_index();
uint32_t user_id = uint32_t(object_index);
uint32_t primitive_offset = 0;
int currIndex = instance_index++;
assert(user_id < blas_lookup.size());
blas_lookup[user_id] = accel_struct_index;
if (geom->geometry_type == Geometry::HAIR) {
/* Build BLAS for curve primitives. */
Hair *const hair = static_cast<Hair *const>(const_cast<Geometry *>(geom));
primitive_offset = uint32_t(hair->curve_segment_offset);
}
else if (geom->geometry_type == Geometry::MESH || geom->geometry_type == Geometry::VOLUME) {
/* Build BLAS for triangle primitives. */
Mesh *const mesh = static_cast<Mesh *const>(const_cast<Geometry *>(geom));
primitive_offset = uint32_t(mesh->prim_offset);
}
else if (geom->geometry_type == Geometry::POINTCLOUD) {
/* Build BLAS for points primitives. */
PointCloud *const pointcloud = static_cast<PointCloud *const>(
const_cast<Geometry *>(geom));
primitive_offset = uint32_t(pointcloud->prim_offset);
}
/* Bake into the appropriate descriptor */
if (motion_blur) {
@@ -871,7 +1010,7 @@ bool BVHMetal::build_TLAS(Progress &progress,
MTLAccelerationStructureMotionInstanceDescriptor &desc = instances[currIndex];
desc.accelerationStructureIndex = accel_struct_index;
desc.userID = user_id;
desc.userID = primitive_offset;
desc.mask = mask;
desc.motionStartTime = 0.0f;
desc.motionEndTime = 1.0f;
@@ -917,9 +1056,10 @@ bool BVHMetal::build_TLAS(Progress &progress,
MTLAccelerationStructureUserIDInstanceDescriptor &desc = instances[currIndex];
desc.accelerationStructureIndex = accel_struct_index;
desc.userID = user_id;
desc.userID = primitive_offset;
desc.mask = mask;
desc.intersectionFunctionTableOffset = 0;
desc.options = MTLAccelerationStructureInstanceOptionOpaque;
float *t = (float *)&desc.transformationMatrix;
if (ob->get_geometry()->is_instanced()) {
@@ -959,6 +1099,7 @@ bool BVHMetal::build_TLAS(Progress &progress,
accelDesc.motionTransformCount = num_motion_transforms;
}
accelDesc.usage |= MTLAccelerationStructureUsageExtendedLimits;
if (!use_fast_trace_bvh) {
accelDesc.usage |= (MTLAccelerationStructureUsageRefit |
MTLAccelerationStructureUsagePreferFastBuild);
@@ -1001,11 +1142,13 @@ bool BVHMetal::build_TLAS(Progress &progress,
/* Cache top and bottom-level acceleration structs */
accel_struct = accel;
blas_array.clear();
blas_array.reserve(all_blas.count);
for (id<MTLAccelerationStructure> blas in all_blas) {
blas_array.push_back(blas);
}
unique_blas_array.clear();
unique_blas_array.reserve(all_blas.count);
[all_blas enumerateObjectsUsingBlock:^(
id<MTLAccelerationStructure> blas, NSUInteger, BOOL *) {
unique_blas_array.push_back(blas);
}];
return true;
}

11
intern/cycles/device/metal/device.mm
View File

@@ -62,12 +62,17 @@ void device_metal_info(vector<DeviceInfo> &devices)
info.has_light_tree = vendor != METAL_GPU_AMD;
info.has_mnee = vendor != METAL_GPU_AMD;
info.use_hardware_raytracing = vendor != METAL_GPU_INTEL;
if (info.use_hardware_raytracing) {
if (@available(macos 11.0, *)) {
info.use_hardware_raytracing = false;
/* MetalRT now uses features exposed in Xcode versions corresponding to macOS 14+, so don't
* expose it in builds from older Xcode versions. */
# if defined(MAC_OS_VERSION_14_0)
if (vendor != METAL_GPU_INTEL) {
if (@available(macos 14.0, *)) {
info.use_hardware_raytracing = device.supportsRaytracing;
}
}
# endif
devices.push_back(info);
device_index++;

1
intern/cycles/device/metal/device_impl.h
View File

@@ -82,7 +82,6 @@ class MetalDevice : public Device {
/* BLAS encoding & lookup */
id<MTLArgumentEncoder> mtlBlasArgEncoder = nil;
id<MTLBuffer> blas_buffer = nil;
id<MTLBuffer> blas_lookup_buffer = nil;
bool use_metalrt = false;
MetalPipelineType kernel_specialization_level = PSO_GENERIC;

29
intern/cycles/device/metal/device_impl.mm
View File

@@ -81,7 +81,7 @@ MetalDevice::MetalDevice(const DeviceInfo &info, Stats &stats, Profiler &profile
mtlDevice = usable_devices[mtlDevId];
device_vendor = MetalInfo::get_device_vendor(mtlDevice);
assert(device_vendor != METAL_GPU_UNKNOWN);
metal_printf("Creating new Cycles device for Metal: %s\n", info.description.c_str());
metal_printf("Creating new Cycles Metal device: %s\n", info.description.c_str());
/* determine default storage mode based on whether UMA is supported */
@@ -549,9 +549,14 @@ void MetalDevice::compile_and_load(int device_id, MetalPipelineType pso_type)
# endif
options.fastMathEnabled = YES;
if (@available(macOS 12.0, *)) {
if (@available(macos 12.0, *)) {
options.languageVersion = MTLLanguageVersion2_4;
}
# if defined(MAC_OS_VERSION_14_0)
if (@available(macos 14.0, *)) {
options.languageVersion = MTLLanguageVersion3_1;
}
# endif
if (getenv("CYCLES_METAL_PROFILING") || getenv("CYCLES_METAL_DEBUG")) {
path_write_text(path_cache_get(string_printf("%s.metal", kernel_type_as_string(pso_type))),
@@ -1372,24 +1377,14 @@ void MetalDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
stats.mem_alloc(blas_buffer.allocatedSize);
for (uint64_t i = 0; i < count; ++i) {
[mtlBlasArgEncoder setArgumentBuffer:blas_buffer
offset:i * mtlBlasArgEncoder.encodedLength];
[mtlBlasArgEncoder setAccelerationStructure:bvhMetalRT->blas_array[i] atIndex:0];
if (bvhMetalRT->blas_array[i]) {
[mtlBlasArgEncoder setArgumentBuffer:blas_buffer
offset:i * mtlBlasArgEncoder.encodedLength];
[mtlBlasArgEncoder setAccelerationStructure:bvhMetalRT->blas_array[i] atIndex:0];
}
}
count = bvhMetalRT->blas_lookup.size();
bufferSize = sizeof(uint32_t) * count;
blas_lookup_buffer = [mtlDevice newBufferWithLength:bufferSize
options:default_storage_mode];
stats.mem_alloc(blas_lookup_buffer.allocatedSize);
memcpy([blas_lookup_buffer contents],
bvhMetalRT -> blas_lookup.data(),
blas_lookup_buffer.allocatedSize);
if (default_storage_mode == MTLResourceStorageModeManaged) {
[blas_buffer didModifyRange:NSMakeRange(0, blas_buffer.length)];
[blas_lookup_buffer didModifyRange:NSMakeRange(0, blas_lookup_buffer.length)];
}
}
}

6
intern/cycles/device/metal/kernel.h
View File

@@ -22,10 +22,8 @@ enum {
METALRT_FUNC_LOCAL_BOX,
METALRT_FUNC_LOCAL_TRI_PRIM,
METALRT_FUNC_LOCAL_BOX_PRIM,
METALRT_FUNC_CURVE_RIBBON,
METALRT_FUNC_CURVE_RIBBON_SHADOW,
METALRT_FUNC_CURVE_ALL,
METALRT_FUNC_CURVE_ALL_SHADOW,
METALRT_FUNC_CURVE,
METALRT_FUNC_CURVE_SHADOW,
METALRT_FUNC_POINT,
METALRT_FUNC_POINT_SHADOW,
METALRT_FUNC_NUM

23
intern/cycles/device/metal/kernel.mm
View File

@@ -493,10 +493,8 @@ void MetalKernelPipeline::compile()
"__anyhit__cycles_metalrt_local_hit_box",
"__anyhit__cycles_metalrt_local_hit_tri_prim",
"__anyhit__cycles_metalrt_local_hit_box_prim",
"__intersection__curve_ribbon",
"__intersection__curve_ribbon_shadow",
"__intersection__curve_all",
"__intersection__curve_all_shadow",
"__intersection__curve",
"__intersection__curve_shadow",
"__intersection__point",
"__intersection__point_shadow",
};
@@ -540,17 +538,8 @@ void MetalKernelPipeline::compile()
id<MTLFunction> point_intersect_default = nil;
id<MTLFunction> point_intersect_shadow = nil;
if (kernel_features & KERNEL_FEATURE_HAIR) {
/* Add curve intersection programs. */
if (kernel_features & KERNEL_FEATURE_HAIR_THICK) {
/* Slower programs for thick hair since that also slows down ribbons.
* Ideally this should not be needed. */
curve_intersect_default = rt_intersection_function[METALRT_FUNC_CURVE_ALL];
curve_intersect_shadow = rt_intersection_function[METALRT_FUNC_CURVE_ALL_SHADOW];
}
else {
curve_intersect_default = rt_intersection_function[METALRT_FUNC_CURVE_RIBBON];
curve_intersect_shadow = rt_intersection_function[METALRT_FUNC_CURVE_RIBBON_SHADOW];
}
curve_intersect_default = rt_intersection_function[METALRT_FUNC_CURVE];
curve_intersect_shadow = rt_intersection_function[METALRT_FUNC_CURVE_SHADOW];
}
if (kernel_features & KERNEL_FEATURE_POINTCLOUD) {
point_intersect_default = rt_intersection_function[METALRT_FUNC_POINT];
@@ -585,8 +574,8 @@ void MetalKernelPipeline::compile()
rt_intersection_function[METALRT_FUNC_LOCAL_BOX_PRIM],
nil];
NSMutableSet *unique_functions = [NSMutableSet
setWithArray:table_functions[METALRT_TABLE_DEFAULT]];
NSMutableSet *unique_functions = [[NSMutableSet alloc] init];
[unique_functions addObjectsFromArray:table_functions[METALRT_TABLE_DEFAULT]];
[unique_functions addObjectsFromArray:table_functions[METALRT_TABLE_SHADOW]];
[unique_functions addObjectsFromArray:table_functions[METALRT_TABLE_LOCAL]];
[unique_functions addObjectsFromArray:table_functions[METALRT_TABLE_LOCAL_PRIM]];

9
intern/cycles/device/metal/queue.mm
View File

@@ -490,9 +490,6 @@ bool MetalDeviceQueue::enqueue(DeviceKernel kernel,
[metal_device_->mtlAncillaryArgEncoder setBuffer:metal_device_->blas_buffer
offset:0
atIndex:8];
[metal_device_->mtlAncillaryArgEncoder setBuffer:metal_device_->blas_lookup_buffer
offset:0
atIndex:9];
}
for (int table = 0; table < METALRT_TABLE_NUM; table++) {
@@ -546,10 +543,8 @@ bool MetalDeviceQueue::enqueue(DeviceKernel kernel,
[mtlComputeCommandEncoder useResource:bvhMetalRT->accel_struct usage:MTLResourceUsageRead];
[mtlComputeCommandEncoder useResource:metal_device_->blas_buffer
usage:MTLResourceUsageRead];
[mtlComputeCommandEncoder useResource:metal_device_->blas_lookup_buffer
usage:MTLResourceUsageRead];
[mtlComputeCommandEncoder useResources:bvhMetalRT->blas_array.data()
count:bvhMetalRT->blas_array.size()
[mtlComputeCommandEncoder useResources:bvhMetalRT->unique_blas_array.data()
count:bvhMetalRT->unique_blas_array.size()
usage:MTLResourceUsageRead];
}
}

342
intern/cycles/kernel/device/metal/bvh.h
View File

@@ -16,12 +16,6 @@ CCL_NAMESPACE_BEGIN
struct MetalRTIntersectionPayload {
RaySelfPrimitives self;
uint visibility;
float u, v;
int prim;
int type;
#if defined(__METALRT_MOTION__)
float time;
#endif
};
struct MetalRTIntersectionLocalPayload {
@@ -37,9 +31,6 @@ struct MetalRTIntersectionLocalPayload {
struct MetalRTIntersectionShadowPayload {
RaySelfPrimitives self;
uint visibility;
#if defined(__METALRT_MOTION__)
float time;
#endif
int state;
float throughput;
short max_hits;
@@ -48,6 +39,98 @@ struct MetalRTIntersectionShadowPayload {
bool result;
};
ccl_device_forceinline bool curve_ribbon_accept(
KernelGlobals kg, float u, float t, ccl_private const Ray *ray, int object, int prim, int type)
{
KernelCurve kcurve = kernel_data_fetch(curves, prim);
int k0 = kcurve.first_key + PRIMITIVE_UNPACK_SEGMENT(type);
int k1 = k0 + 1;
int ka = max(k0 - 1, kcurve.first_key);
int kb = min(k1 + 1, kcurve.first_key + kcurve.num_keys - 1);
/* We can ignore motion blur here because we don't need the positions, and it doesn't affect the
* radius. */
float radius[4];
radius[0] = kernel_data_fetch(curve_keys, ka).w;
radius[1] = kernel_data_fetch(curve_keys, k0).w;
radius[2] = kernel_data_fetch(curve_keys, k1).w;
radius[3] = kernel_data_fetch(curve_keys, kb).w;
const float r = metal::catmull_rom(u, radius[0], radius[1], radius[2], radius[3]);
/* MPJ TODO: Can we ignore motion and/or object transforms here? Depends on scaling? */
float3 ray_P = ray->P;
float3 ray_D = ray->D;
if (!(kernel_data_fetch(object_flag, object) & SD_OBJECT_TRANSFORM_APPLIED)) {
float3 idir;
#if defined(__METALRT_MOTION__)
bvh_instance_motion_push(NULL, object, ray, &ray_P, &ray_D, &idir);
#else
bvh_instance_push(NULL, object, ray, &ray_P, &ray_D, &idir);
#endif
}
/* ignore self intersections */
const float avoidance_factor = 2.0f;
return t * len(ray_D) > avoidance_factor * r;
}
ccl_device_forceinline float curve_ribbon_v(
KernelGlobals kg, float u, float t, ccl_private const Ray *ray, int object, int prim, int type)
{
#if defined(__METALRT_MOTION__)
float time = ray->time;
#else
float time = 0.0f;
#endif
const bool is_motion = (type & PRIMITIVE_MOTION);
KernelCurve kcurve = kernel_data_fetch(curves, prim);
int k0 = kcurve.first_key + PRIMITIVE_UNPACK_SEGMENT(type);
int k1 = k0 + 1;
int ka = max(k0 - 1, kcurve.first_key);
int kb = min(k1 + 1, kcurve.first_key + kcurve.num_keys - 1);
float4 curve[4];
if (!is_motion) {
curve[0] = kernel_data_fetch(curve_keys, ka);
curve[1] = kernel_data_fetch(curve_keys, k0);
curve[2] = kernel_data_fetch(curve_keys, k1);
curve[3] = kernel_data_fetch(curve_keys, kb);
}
else {
motion_curve_keys(kg, object, prim, time, ka, k0, k1, kb, curve);
}
float3 ray_P = ray->P;
float3 ray_D = ray->D;
if (!(kernel_data_fetch(object_flag, object) & SD_OBJECT_TRANSFORM_APPLIED)) {
float3 idir;
#if defined(__METALRT_MOTION__)
bvh_instance_motion_push(NULL, object, ray, &ray_P, &ray_D, &idir);
#else
bvh_instance_push(NULL, object, ray, &ray_P, &ray_D, &idir);
#endif
}
const float4 P_curve4 = metal::catmull_rom(u, curve[0], curve[1], curve[2], curve[3]);
const float r_curve = P_curve4.w;
float3 P = ray_P + ray_D * t;
const float3 P_curve = float4_to_float3(P_curve4);
const float4 dPdu4 = metal::catmull_rom_derivative(u, curve[0], curve[1], curve[2], curve[3]);
const float3 dPdu = float4_to_float3(dPdu4);
const float3 tangent = normalize(dPdu);
const float3 bitangent = normalize(cross(tangent, -ray_D));
float v = dot(P - P_curve, bitangent) / r_curve;
return clamp(v, -1.0, 1.0f);
}
/* Scene intersection. */
ccl_device_intersect bool scene_intersect(KernelGlobals kg,
@@ -79,41 +162,34 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg,
metal::raytracing::ray r(ray->P, ray->D, ray->tmin, ray->tmax);
metalrt_intersector_type metalrt_intersect;
metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque);
metalrt_intersect.assume_geometry_type(
metal::raytracing::geometry_type::triangle |
(kernel_data.bvh.have_curves ? metal::raytracing::geometry_type::curve :
metal::raytracing::geometry_type::none) |
(kernel_data.bvh.have_points ? metal::raytracing::geometry_type::bounding_box :
metal::raytracing::geometry_type::none));
bool triangle_only = !kernel_data.bvh.have_curves && !kernel_data.bvh.have_points;
if (triangle_only) {
metalrt_intersect.assume_geometry_type(metal::raytracing::geometry_type::triangle);
if (visibility & PATH_RAY_SHADOW_OPAQUE) {
metalrt_intersect.accept_any_intersection(true);
}
MetalRTIntersectionPayload payload;
payload.self = ray->self;
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 behavior is any-hit. */
}
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,
visibility,
ray->time,
metal_ancillaries->ift_default,
payload);
#else
intersection = metalrt_intersect.intersect(
r, metal_ancillaries->accel_struct, ray_mask, metal_ancillaries->ift_default, payload);
r, metal_ancillaries->accel_struct, visibility, metal_ancillaries->ift_default, payload);
#endif
if (intersection.type == intersection_type::none) {
@@ -123,23 +199,71 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg,
return false;
}
isect->t = intersection.distance;
isect->prim = payload.prim;
isect->type = payload.type;
isect->object = intersection.user_instance_id;
isect->object = intersection.instance_id;
isect->t = intersection.distance;
if (intersection.type == intersection_type::triangle) {
isect->prim = intersection.primitive_id + intersection.user_instance_id;
isect->type = kernel_data_fetch(objects, intersection.instance_id).primitive_type;
isect->u = intersection.triangle_barycentric_coord.x;
isect->v = intersection.triangle_barycentric_coord.y;
}
else {
isect->u = payload.u;
isect->v = payload.v;
else if (kernel_data.bvh.have_curves && intersection.type == intersection_type::curve) {
int prim = intersection.primitive_id + intersection.user_instance_id;
const KernelCurveSegment segment = kernel_data_fetch(curve_segments, prim);
isect->prim = segment.prim;
isect->type = segment.type;
isect->u = intersection.curve_parameter;
if (segment.type & PRIMITIVE_CURVE_RIBBON) {
isect->v = curve_ribbon_v(kg,
intersection.curve_parameter,
intersection.distance,
ray,
intersection.instance_id,
segment.prim,
segment.type);
}
else {
isect->v = 0.0f;
}
}
else if (kernel_data.bvh.have_points && intersection.type == intersection_type::bounding_box) {
const int object = intersection.instance_id;
const uint prim = intersection.primitive_id + intersection.user_instance_id;
const int prim_type = kernel_data_fetch(objects, object).primitive_type;
if (!(kernel_data_fetch(object_flag, object) & SD_OBJECT_TRANSFORM_APPLIED)) {
float3 idir;
#if defined(__METALRT_MOTION__)
bvh_instance_motion_push(NULL, object, ray, &r.origin, &r.direction, &idir);
#else
bvh_instance_push(NULL, object, ray, &r.origin, &r.direction, &idir);
#endif
}
if (prim_type & PRIMITIVE_POINT) {
if (!point_intersect(NULL,
isect,
r.origin,
r.direction,
ray->tmin,
ray->tmax,
object,
prim,
ray->time,
prim_type))
{
/* Shouldn't get here */
kernel_assert(!"Intersection mismatch");
isect->t = ray->tmax;
isect->type = PRIMITIVE_NONE;
return false;
}
return true;
}
}
return isect->type != PRIMITIVE_NONE;
return true;
}
#ifdef __BVH_LOCAL__
@@ -198,25 +322,18 @@ ccl_device_intersect bool scene_intersect_local(KernelGlobals kg,
# if defined(__METALRT_MOTION__)
metalrt_intersector_type metalrt_intersect;
typename metalrt_intersector_type::result_type intersection;
metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque);
bool triangle_only = !kernel_data.bvh.have_curves && !kernel_data.bvh.have_points;
if (triangle_only) {
metalrt_intersect.assume_geometry_type(metal::raytracing::geometry_type::triangle);
}
intersection = metalrt_intersect.intersect(
r, metal_ancillaries->accel_struct, 0xFF, ray->time, metal_ancillaries->ift_local, payload);
# else
metalrt_blas_intersector_type metalrt_intersect;
typename metalrt_blas_intersector_type::result_type intersection;
# endif
metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque);
bool triangle_only = !kernel_data.bvh.have_curves && !kernel_data.bvh.have_points;
if (triangle_only) {
metalrt_intersect.assume_geometry_type(metal::raytracing::geometry_type::triangle);
}
metalrt_intersect.assume_geometry_type(
metal::raytracing::geometry_type::triangle |
(kernel_data.bvh.have_curves ? metal::raytracing::geometry_type::curve :
metal::raytracing::geometry_type::none) |
(kernel_data.bvh.have_points ? metal::raytracing::geometry_type::bounding_box :
metal::raytracing::geometry_type::none));
// if we know we are going to get max one hit, like for random-sss-walk we can
// optimize and accept the first hit
@@ -224,8 +341,10 @@ ccl_device_intersect bool scene_intersect_local(KernelGlobals kg,
metalrt_intersect.accept_any_intersection(true);
}
int blas_index = metal_ancillaries->blas_userID_to_index_lookUp[local_object];
# if defined(__METALRT_MOTION__)
intersection = metalrt_intersect.intersect(
r, metal_ancillaries->accel_struct, ~0, ray->time, metal_ancillaries->ift_local, payload);
# else
if (!(kernel_data_fetch(object_flag, local_object) & SD_OBJECT_TRANSFORM_APPLIED)) {
// transform the ray into object's local space
Transform itfm = kernel_data_fetch(objects, local_object).itfm;
@@ -235,7 +354,7 @@ ccl_device_intersect bool scene_intersect_local(KernelGlobals kg,
intersection = metalrt_intersect.intersect(
r,
metal_ancillaries->blas_accel_structs[blas_index].blas,
metal_ancillaries->blas_accel_structs[local_object].blas,
metal_ancillaries->ift_local_prim,
payload);
# endif
@@ -278,13 +397,13 @@ ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg,
metal::raytracing::ray r(ray->P, ray->D, ray->tmin, ray->tmax);
metalrt_intersector_type metalrt_intersect;
metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque);
bool triangle_only = !kernel_data.bvh.have_curves && !kernel_data.bvh.have_points;
if (triangle_only) {
metalrt_intersect.assume_geometry_type(metal::raytracing::geometry_type::triangle);
}
metalrt_intersect.assume_geometry_type(
metal::raytracing::geometry_type::triangle |
(kernel_data.bvh.have_curves ? metal::raytracing::geometry_type::curve :
metal::raytracing::geometry_type::none) |
(kernel_data.bvh.have_points ? metal::raytracing::geometry_type::bounding_box :
metal::raytracing::geometry_type::none));
MetalRTIntersectionShadowPayload payload;
payload.self = ray->self;
@@ -296,24 +415,18 @@ ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg,
payload.result = false;
payload.state = state;
uint ray_mask = visibility & 0xFF;
if (0 == ray_mask && (visibility & ~0xFF) != 0) {
ray_mask = 0xFF;
}
typename metalrt_intersector_type::result_type intersection;
# if defined(__METALRT_MOTION__)
payload.time = ray->time;
intersection = metalrt_intersect.intersect(r,
metal_ancillaries->accel_struct,
ray_mask,
visibility,
ray->time,
metal_ancillaries->ift_shadow,
payload);
# else
intersection = metalrt_intersect.intersect(
r, metal_ancillaries->accel_struct, ray_mask, metal_ancillaries->ift_shadow, payload);
r, metal_ancillaries->accel_struct, visibility, metal_ancillaries->ift_shadow, payload);
# endif
*num_recorded_hits = payload.num_recorded_hits;
@@ -347,13 +460,13 @@ ccl_device_intersect bool scene_intersect_volume(KernelGlobals kg,
metal::raytracing::ray r(ray->P, ray->D, ray->tmin, ray->tmax);
metalrt_intersector_type metalrt_intersect;
metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque);
bool triangle_only = !kernel_data.bvh.have_curves && !kernel_data.bvh.have_points;
if (triangle_only) {
metalrt_intersect.assume_geometry_type(metal::raytracing::geometry_type::triangle);
}
metalrt_intersect.assume_geometry_type(
metal::raytracing::geometry_type::triangle |
(kernel_data.bvh.have_curves ? metal::raytracing::geometry_type::curve :
metal::raytracing::geometry_type::none) |
(kernel_data.bvh.have_points ? metal::raytracing::geometry_type::bounding_box :
metal::raytracing::geometry_type::none));
MetalRTIntersectionPayload payload;
payload.self = ray->self;
@@ -361,43 +474,86 @@ ccl_device_intersect bool scene_intersect_volume(KernelGlobals kg,
typename metalrt_intersector_type::result_type intersection;
uint ray_mask = visibility & 0xFF;
if (0 == ray_mask && (visibility & ~0xFF) != 0) {
ray_mask = 0xFF;
}
# if defined(__METALRT_MOTION__)
payload.time = ray->time;
intersection = metalrt_intersect.intersect(r,
metal_ancillaries->accel_struct,
ray_mask,
visibility,
ray->time,
metal_ancillaries->ift_default,
payload);
# else
intersection = metalrt_intersect.intersect(
r, metal_ancillaries->accel_struct, ray_mask, metal_ancillaries->ift_default, payload);
r, metal_ancillaries->accel_struct, visibility, metal_ancillaries->ift_default, payload);
# endif
if (intersection.type == intersection_type::none) {
return false;
}
isect->prim = payload.prim;
isect->type = payload.type;
isect->object = intersection.user_instance_id;
isect->t = intersection.distance;
if (intersection.type == intersection_type::triangle) {
else if (intersection.type == intersection_type::triangle) {
isect->prim = intersection.primitive_id + intersection.user_instance_id;
isect->type = kernel_data_fetch(objects, intersection.instance_id).primitive_type;
isect->u = intersection.triangle_barycentric_coord.x;
isect->v = intersection.triangle_barycentric_coord.y;
isect->object = intersection.instance_id;
isect->t = intersection.distance;
}
else {
isect->u = payload.u;
isect->v = payload.v;
else if (kernel_data.bvh.have_curves && intersection.type == intersection_type::curve) {
int prim = intersection.primitive_id + intersection.user_instance_id;
const KernelCurveSegment segment = kernel_data_fetch(curve_segments, prim);
isect->prim = segment.prim;
isect->type = segment.type;
isect->u = intersection.curve_parameter;
if (segment.type & PRIMITIVE_CURVE_RIBBON) {
isect->v = curve_ribbon_v(kg,
intersection.curve_parameter,
intersection.distance,
ray,
intersection.instance_id,
segment.prim,
segment.type);
}
else {
isect->v = 0.0f;
}
}
else if (kernel_data.bvh.have_points && intersection.type == intersection_type::bounding_box) {
const int object = intersection.instance_id;
const uint prim = intersection.primitive_id + intersection.user_instance_id;
const int prim_type = kernel_data_fetch(objects, intersection.instance_id).primitive_type;
isect->object = object;
if (!(kernel_data_fetch(object_flag, object) & SD_OBJECT_TRANSFORM_APPLIED)) {
float3 idir;
# if defined(__METALRT_MOTION__)
bvh_instance_motion_push(NULL, object, ray, &r.origin, &r.direction, &idir);
# else
bvh_instance_push(NULL, object, ray, &r.origin, &r.direction, &idir);
# endif
}
if (prim_type & PRIMITIVE_POINT) {
if (!point_intersect(NULL,
isect,
r.origin,
r.direction,
ray->tmin,
ray->tmax,
intersection.instance_id,
prim,
ray->time,
prim_type))
{
/* Shouldn't get here */
kernel_assert(!"Intersection mismatch");
return false;
}
return true;
}
}
return isect->type != PRIMITIVE_NONE;
return true;
}
#endif

23
intern/cycles/kernel/device/metal/compat.h
View File

@@ -27,6 +27,8 @@ using namespace metal::raytracing;
#pragma clang diagnostic ignored "-Wunused-variable"
#pragma clang diagnostic ignored "-Wsign-compare"
#pragma clang diagnostic ignored "-Wuninitialized"
#pragma clang diagnostic ignored "-Wc++17-extensions"
#pragma clang diagnostic ignored "-Wmacro-redefined"
/* Qualifiers */
@@ -280,17 +282,23 @@ ccl_device_forceinline uchar4 make_uchar4(const uchar x,
# endif /* __METALRT_MOTION__ */
typedef acceleration_structure<METALRT_TAGS> metalrt_as_type;
typedef intersection_function_table<triangle_data, METALRT_TAGS> metalrt_ift_type;
typedef metal::raytracing::intersector<triangle_data, METALRT_TAGS> metalrt_intersector_type;
typedef intersection_function_table<triangle_data, curve_data, METALRT_TAGS, extended_limits>
metalrt_ift_type;
typedef metal::raytracing::intersector<triangle_data, curve_data, METALRT_TAGS, extended_limits>
metalrt_intersector_type;
# if defined(__METALRT_MOTION__)
typedef acceleration_structure<primitive_motion> metalrt_blas_as_type;
typedef intersection_function_table<triangle_data, primitive_motion> metalrt_blas_ift_type;
typedef metal::raytracing::intersector<triangle_data, primitive_motion>
metalrt_blas_intersector_type;
typedef intersection_function_table<triangle_data, curve_data, primitive_motion, extended_limits>
metalrt_blas_ift_type;
typedef metal::raytracing::
intersector<triangle_data, curve_data, primitive_motion, extended_limits>
metalrt_blas_intersector_type;
# else
typedef acceleration_structure<> metalrt_blas_as_type;
typedef intersection_function_table<triangle_data> metalrt_blas_ift_type;
typedef metal::raytracing::intersector<triangle_data> metalrt_blas_intersector_type;
typedef intersection_function_table<triangle_data, curve_data, extended_limits>
metalrt_blas_ift_type;
typedef metal::raytracing::intersector<triangle_data, curve_data, extended_limits>
metalrt_blas_intersector_type;
# endif
#endif /* __METALRT__ */
@@ -326,7 +334,6 @@ struct MetalAncillaries {
metalrt_ift_type ift_local;
metalrt_blas_ift_type ift_local_prim;
constant MetalRTBlasWrapper *blas_accel_structs;
constant int *blas_userID_to_index_lookUp;
#endif
};

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

@@ -26,13 +26,13 @@ struct BoundingBoxIntersectionResult {
float distance [[distance]];
};
/* For a triangle intersection function. */
struct TriangleIntersectionResult {
/* For a primitive intersection function. */
struct PrimitiveIntersectionResult {
bool accept [[accept_intersection]];
bool continue_search [[continue_search]];
};
enum { METALRT_HIT_TRIANGLE, METALRT_HIT_BOUNDING_BOX };
enum { METALRT_HIT_TRIANGLE, METALRT_HIT_CURVE, METALRT_HIT_BOUNDING_BOX };
/* Hit functions. */
@@ -40,20 +40,17 @@ template<typename TReturn, uint intersection_type>
TReturn metalrt_local_hit(constant KernelParamsMetal &launch_params_metal,
ray_data MetalKernelContext::MetalRTIntersectionLocalPayload &payload,
const uint object,
const uint primitive_id,
const uint prim,
const float2 barycentrics,
const float ray_tmax)
{
TReturn result;
# ifdef __BVH_LOCAL__
uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object);
#ifdef __BVH_LOCAL__
MetalKernelContext context(launch_params_metal);
if ((object != payload.local_object) || context.intersection_skip_self_local(payload.self, prim))
{
/* Only intersect with matching object and skip self-intersection. */
if ((object != payload.local_object) || context.intersection_skip_self_local(payload.self, prim)) {
/* Only intersect with matching object and skip self-intersecton. */
result.accept = false;
result.continue_search = true;
return result;
@@ -124,7 +121,7 @@ TReturn metalrt_local_hit(constant KernelParamsMetal &launch_params_metal,
# endif
}
[[intersection(triangle, triangle_data)]] TriangleIntersectionResult
[[intersection(triangle, triangle_data, curve_data)]] PrimitiveIntersectionResult
__anyhit__cycles_metalrt_local_hit_tri_prim(
constant KernelParamsMetal &launch_params_metal [[buffer(1)]],
ray_data MetalKernelContext::MetalRTIntersectionLocalPayload &payload [[payload]],
@@ -132,27 +129,30 @@ __anyhit__cycles_metalrt_local_hit_tri_prim(
float2 barycentrics [[barycentric_coord]],
float ray_tmax [[distance]])
{
uint prim = primitive_id + kernel_data_fetch(object_prim_offset, payload.local_object);
/* instance_id, aka the user_id has been removed. If we take this function we optimized the
* SSS for starting traversal from a primitive acceleration structure instead of the root of the
* global AS. this means we will always be intersecting the correct object no need for the
* user-id to check */
return metalrt_local_hit<TriangleIntersectionResult, METALRT_HIT_TRIANGLE>(
launch_params_metal, payload, payload.local_object, primitive_id, barycentrics, ray_tmax);
return metalrt_local_hit<PrimitiveIntersectionResult, METALRT_HIT_TRIANGLE>(
launch_params_metal, payload, payload.local_object, prim, barycentrics, ray_tmax);
}
[[intersection(triangle, triangle_data, METALRT_TAGS)]] TriangleIntersectionResult
[[intersection(triangle, triangle_data, curve_data, METALRT_TAGS, extended_limits)]] PrimitiveIntersectionResult
__anyhit__cycles_metalrt_local_hit_tri(
constant KernelParamsMetal &launch_params_metal [[buffer(1)]],
ray_data MetalKernelContext::MetalRTIntersectionLocalPayload &payload [[payload]],
uint instance_id [[user_instance_id]],
uint instance_id [[instance_id]],
uint primitive_id [[primitive_id]],
uint primitive_id_offset [[user_instance_id]],
float2 barycentrics [[barycentric_coord]],
float ray_tmax [[distance]])
{
return metalrt_local_hit<TriangleIntersectionResult, METALRT_HIT_TRIANGLE>(
launch_params_metal, payload, instance_id, primitive_id, barycentrics, ray_tmax);
return metalrt_local_hit<PrimitiveIntersectionResult, METALRT_HIT_TRIANGLE>(
launch_params_metal, payload, instance_id, primitive_id + primitive_id_offset, barycentrics, ray_tmax);
}
[[intersection(bounding_box, triangle_data, METALRT_TAGS)]] BoundingBoxIntersectionResult
[[intersection(bounding_box, triangle_data, curve_data, METALRT_TAGS, extended_limits)]] BoundingBoxIntersectionResult
__anyhit__cycles_metalrt_local_hit_box(const float ray_tmax [[max_distance]])
{
/* unused function */
@@ -163,7 +163,7 @@ __anyhit__cycles_metalrt_local_hit_box(const float ray_tmax [[max_distance]])
return result;
}
[[intersection(bounding_box, triangle_data)]] BoundingBoxIntersectionResult
[[intersection(bounding_box, triangle_data, curve_data)]] BoundingBoxIntersectionResult
__anyhit__cycles_metalrt_local_hit_box_prim(const float ray_tmax [[max_distance]])
{
/* unused function */
@@ -180,30 +180,32 @@ bool metalrt_shadow_all_hit(constant KernelParamsMetal &launch_params_metal,
uint object,
uint prim,
const float2 barycentrics,
const float ray_tmax)
const float ray_tmax,
const float t = 0.0f,
ccl_private const Ray *ray = NULL
)
{
# ifdef __SHADOW_RECORD_ALL__
# ifdef __VISIBILITY_FLAG__
const uint visibility = payload.visibility;
if ((kernel_data_fetch(objects, object).visibility & visibility) == 0) {
/* continue search */
return true;
}
# endif
const float u = barycentrics.x;
const float v = barycentrics.y;
#ifdef __SHADOW_RECORD_ALL__
float u = barycentrics.x;
float v = barycentrics.y;
const int prim_type = kernel_data_fetch(objects, object).primitive_type;
int type = prim_type;
# ifdef __HAIR__
if (intersection_type != METALRT_HIT_TRIANGLE) {
if ((prim_type == PRIMITIVE_CURVE_THICK || prim_type == PRIMITIVE_CURVE_RIBBON)) {
const KernelCurveSegment segment = kernel_data_fetch(curve_segments, prim);
type = segment.type;
prim = segment.prim;
int type;
/* Filter out curve end-caps. */
if (u == 0.0f || u == 1.0f) {
# ifdef __HAIR__
if constexpr (intersection_type == METALRT_HIT_CURVE) {
const KernelCurveSegment segment = kernel_data_fetch(curve_segments, prim);
type = segment.type;
prim = segment.prim;
/* Filter out curve end-caps. */
if (u == 0.0f || u == 1.0f) {
/* continue search */
return true;
}
if (type & PRIMITIVE_CURVE_RIBBON) {
MetalKernelContext context(launch_params_metal);
if (!context.curve_ribbon_accept(NULL, u, t, ray, object, prim, type)) {
/* continue search */
return true;
}
@@ -211,6 +213,17 @@ bool metalrt_shadow_all_hit(constant KernelParamsMetal &launch_params_metal,
}
# endif
if constexpr (intersection_type == METALRT_HIT_BOUNDING_BOX) {
/* Point. */
type = kernel_data_fetch(objects, object).primitive_type;
u = 0.0f;
v = 0.0f;
}
if constexpr (intersection_type == METALRT_HIT_TRIANGLE) {
type = prim_type;
}
MetalKernelContext context(launch_params_metal);
if (context.intersection_skip_self_shadow(payload.self, object, prim)) {
@@ -244,8 +257,9 @@ bool metalrt_shadow_all_hit(constant KernelParamsMetal &launch_params_metal,
return false;
}
# ifdef __HAIR__
/* Always use baked shadow transparency for curves. */
if (type & PRIMITIVE_CURVE) {
if constexpr (intersection_type == METALRT_HIT_CURVE) {
float throughput = payload.throughput;
throughput *= context.intersection_curve_shadow_transparency(nullptr, object, prim, type, u);
payload.throughput = throughput;
@@ -260,6 +274,7 @@ bool metalrt_shadow_all_hit(constant KernelParamsMetal &launch_params_metal,
return true;
}
}
# endif
payload.num_hits += 1;
payload.num_recorded_hits += 1;
@@ -305,25 +320,26 @@ bool metalrt_shadow_all_hit(constant KernelParamsMetal &launch_params_metal,
return true;
}
[[intersection(triangle, triangle_data, METALRT_TAGS)]] TriangleIntersectionResult
[[intersection(triangle, triangle_data, curve_data, METALRT_TAGS, extended_limits)]] PrimitiveIntersectionResult
__anyhit__cycles_metalrt_shadow_all_hit_tri(
constant KernelParamsMetal &launch_params_metal [[buffer(1)]],
ray_data MetalKernelContext::MetalRTIntersectionShadowPayload &payload [[payload]],
unsigned int object [[user_instance_id]],
unsigned int primitive_id [[primitive_id]],
float2 barycentrics [[barycentric_coord]],
float ray_tmax [[distance]])
const unsigned int object [[instance_id]],
const unsigned int primitive_id [[primitive_id]],
const uint primitive_id_offset [[user_instance_id]],
const float2 barycentrics [[barycentric_coord]],
const float ray_tmax [[distance]])
{
uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object);
uint prim = primitive_id + primitive_id_offset;
TriangleIntersectionResult result;
PrimitiveIntersectionResult result;
result.continue_search = metalrt_shadow_all_hit<METALRT_HIT_TRIANGLE>(
launch_params_metal, payload, object, prim, barycentrics, ray_tmax);
result.accept = !result.continue_search;
return result;
}
[[intersection(bounding_box, triangle_data, METALRT_TAGS)]] BoundingBoxIntersectionResult
[[intersection(bounding_box, triangle_data, curve_data, METALRT_TAGS, extended_limits)]] BoundingBoxIntersectionResult
__anyhit__cycles_metalrt_shadow_all_hit_box(const float ray_tmax [[max_distance]])
{
/* unused function */
@@ -340,40 +356,38 @@ inline TReturnType metalrt_visibility_test(
ray_data MetalKernelContext::MetalRTIntersectionPayload &payload,
const uint object,
uint prim,
const float u)
const float u,
const float t = 0.0f,
ccl_private const Ray *ray = NULL
)
{
TReturnType result;
# ifdef __HAIR__
const int type = kernel_data_fetch(objects, object).primitive_type;
if (intersection_type == METALRT_HIT_BOUNDING_BOX &&
(type == PRIMITIVE_CURVE_THICK || type == PRIMITIVE_CURVE_RIBBON))
{
#ifdef __HAIR__
if constexpr (intersection_type == METALRT_HIT_CURVE) {
/* Filter out curve end-caps. */
if (u == 0.0f || u == 1.0f) {
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(NULL, u, t, ray, object, prim, type)) {
result.accept = false;
result.continue_search = true;
return result;
}
}
}
# endif
uint visibility = payload.visibility;
# ifdef __VISIBILITY_FLAG__
if ((kernel_data_fetch(objects, object).visibility & visibility) == 0) {
result.accept = false;
result.continue_search = true;
return result;
}
# endif
if (intersection_type == METALRT_HIT_TRIANGLE) {
}
# ifdef __HAIR__
else {
prim = kernel_data_fetch(curve_segments, prim).prim;
}
# endif
MetalKernelContext context(launch_params_metal);
@@ -411,25 +425,22 @@ inline TReturnType metalrt_visibility_test(
return result;
}
[[intersection(triangle, triangle_data, METALRT_TAGS)]] TriangleIntersectionResult
[[intersection(triangle, triangle_data, curve_data, METALRT_TAGS, extended_limits)]] PrimitiveIntersectionResult
__anyhit__cycles_metalrt_visibility_test_tri(
constant KernelParamsMetal &launch_params_metal [[buffer(1)]],
ray_data MetalKernelContext::MetalRTIntersectionPayload &payload [[payload]],
unsigned int object [[user_instance_id]],
unsigned int primitive_id [[primitive_id]])
const unsigned int object [[instance_id]],
const uint primitive_id_offset [[user_instance_id]],
const unsigned int primitive_id [[primitive_id]])
{
uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object);
TriangleIntersectionResult result =
metalrt_visibility_test<TriangleIntersectionResult, METALRT_HIT_TRIANGLE>(
uint prim = primitive_id + primitive_id_offset;
PrimitiveIntersectionResult result =
metalrt_visibility_test<PrimitiveIntersectionResult, METALRT_HIT_TRIANGLE>(
launch_params_metal, payload, object, prim, 0.0f);
if (result.accept) {
payload.prim = prim;
payload.type = kernel_data_fetch(objects, object).primitive_type;
}
return result;
}
[[intersection(bounding_box, triangle_data, METALRT_TAGS)]] BoundingBoxIntersectionResult
[[intersection(bounding_box, triangle_data, curve_data, METALRT_TAGS, extended_limits)]] BoundingBoxIntersectionResult
__anyhit__cycles_metalrt_visibility_test_box(const float ray_tmax [[max_distance]])
{
/* Unused function */
@@ -442,230 +453,72 @@ __anyhit__cycles_metalrt_visibility_test_box(const float ray_tmax [[max_distance
/* Primitive intersection functions. */
# ifdef __HAIR__
ccl_device_inline void metalrt_intersection_curve(
constant KernelParamsMetal &launch_params_metal,
ray_data MetalKernelContext::MetalRTIntersectionPayload &payload,
const uint object,
const uint prim,
const uint type,
const float3 ray_P,
const float3 ray_D,
float time,
const float ray_tmin,
const float ray_tmax,
thread BoundingBoxIntersectionResult &result)
#ifdef __HAIR__
[[intersection(curve, triangle_data, curve_data, METALRT_TAGS, extended_limits)]] PrimitiveIntersectionResult
__intersection__curve(constant KernelParamsMetal &launch_params_metal [[buffer(1)]],
ray_data MetalKernelContext::MetalRTIntersectionPayload &payload
[[payload]],
const uint object [[instance_id]],
const uint primitive_id [[primitive_id]],
const uint primitive_id_offset [[user_instance_id]],
float distance [[distance]],
const float3 ray_P [[origin]],
const float3 ray_D [[direction]],
float u [[curve_parameter]],
const float ray_tmin [[min_distance]],
const float ray_tmax [[max_distance]]
# if defined(__METALRT_MOTION__)
,const float time [[time]]
# endif
)
{
# ifdef __VISIBILITY_FLAG__
const uint visibility = payload.visibility;
if ((kernel_data_fetch(objects, object).visibility & visibility) == 0) {
return;
}
# endif
uint prim = primitive_id + primitive_id_offset;
Intersection isect;
isect.t = ray_tmax;
Ray ray;
ray.P = ray_P;
ray.D = ray_D;
#if defined(__METALRT_MOTION__)
ray.time = time;
#endif
MetalKernelContext context(launch_params_metal);
if (context.curve_intersect(
NULL, &isect, ray_P, ray_D, ray_tmin, isect.t, object, prim, time, type))
{
result = metalrt_visibility_test<BoundingBoxIntersectionResult, METALRT_HIT_BOUNDING_BOX>(
launch_params_metal, payload, object, prim, isect.u);
if (result.accept) {
result.distance = isect.t;
payload.u = isect.u;
payload.v = isect.v;
payload.prim = prim;
payload.type = type;
}
}
}
ccl_device_inline void metalrt_intersection_curve_shadow(
constant KernelParamsMetal &launch_params_metal,
ray_data MetalKernelContext::MetalRTIntersectionShadowPayload &payload,
const uint object,
const uint prim,
const uint type,
const float3 ray_P,
const float3 ray_D,
float time,
const float ray_tmin,
const float ray_tmax,
thread BoundingBoxIntersectionResult &result)
{
# ifdef __VISIBILITY_FLAG__
const uint visibility = payload.visibility;
if ((kernel_data_fetch(objects, object).visibility & visibility) == 0) {
return;
}
# endif
Intersection isect;
isect.t = ray_tmax;
MetalKernelContext context(launch_params_metal);
if (context.curve_intersect(
NULL, &isect, ray_P, ray_D, ray_tmin, isect.t, object, prim, time, type))
{
result.continue_search = metalrt_shadow_all_hit<METALRT_HIT_BOUNDING_BOX>(
launch_params_metal, payload, object, prim, float2(isect.u, isect.v), ray_tmax);
result.accept = !result.continue_search;
}
}
[[intersection(bounding_box, triangle_data, METALRT_TAGS)]] BoundingBoxIntersectionResult
__intersection__curve_ribbon(constant KernelParamsMetal &launch_params_metal [[buffer(1)]],
ray_data MetalKernelContext::MetalRTIntersectionPayload &payload
[[payload]],
const uint object [[user_instance_id]],
const uint primitive_id [[primitive_id]],
const float3 ray_P [[origin]],
const float3 ray_D [[direction]],
const float ray_tmin [[min_distance]],
const float ray_tmax [[max_distance]])
{
uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object);
const KernelCurveSegment segment = kernel_data_fetch(curve_segments, prim);
BoundingBoxIntersectionResult result;
result.accept = false;
result.continue_search = true;
result.distance = ray_tmax;
if (segment.type & PRIMITIVE_CURVE_RIBBON) {
metalrt_intersection_curve(launch_params_metal,
payload,
object,
segment.prim,
segment.type,
ray_P,
ray_D,
# if defined(__METALRT_MOTION__)
payload.time,
# else
0.0f,
# endif
ray_tmin,
ray_tmax,
result);
}
PrimitiveIntersectionResult result =
metalrt_visibility_test<PrimitiveIntersectionResult, METALRT_HIT_CURVE>(
launch_params_metal, payload, object, prim, u, distance, &ray);
return result;
}
[[intersection(bounding_box, triangle_data, METALRT_TAGS)]] BoundingBoxIntersectionResult
__intersection__curve_ribbon_shadow(
[[intersection(curve, triangle_data, curve_data, METALRT_TAGS, extended_limits)]] PrimitiveIntersectionResult
__intersection__curve_shadow(
constant KernelParamsMetal &launch_params_metal [[buffer(1)]],
ray_data MetalKernelContext::MetalRTIntersectionShadowPayload &payload [[payload]],
const uint object [[user_instance_id]],
const uint object [[instance_id]],
const uint primitive_id [[primitive_id]],
const uint primitive_id_offset [[user_instance_id]],
const float3 ray_P [[origin]],
const float3 ray_D [[direction]],
float u [[curve_parameter]],
float t [[distance]],
# if defined(__METALRT_MOTION__)
const float time [[time]],
# endif
const float ray_tmin [[min_distance]],
const float ray_tmax [[max_distance]])
{
uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object);
const KernelCurveSegment segment = kernel_data_fetch(curve_segments, prim);
uint prim = primitive_id + primitive_id_offset;
BoundingBoxIntersectionResult result;
result.accept = false;
result.continue_search = true;
result.distance = ray_tmax;
PrimitiveIntersectionResult result;
if (segment.type & PRIMITIVE_CURVE_RIBBON) {
metalrt_intersection_curve_shadow(launch_params_metal,
payload,
object,
segment.prim,
segment.type,
ray_P,
ray_D,
# if defined(__METALRT_MOTION__)
payload.time,
# else
0.0f,
# endif
ray_tmin,
ray_tmax,
result);
}
Ray ray;
ray.P = ray_P;
ray.D = ray_D;
#if defined(__METALRT_MOTION__)
ray.time = time;
#endif
return result;
}
[[intersection(bounding_box, triangle_data, METALRT_TAGS)]] BoundingBoxIntersectionResult
__intersection__curve_all(constant KernelParamsMetal &launch_params_metal [[buffer(1)]],
ray_data MetalKernelContext::MetalRTIntersectionPayload &payload
[[payload]],
const uint object [[user_instance_id]],
const uint primitive_id [[primitive_id]],
const float3 ray_P [[origin]],
const float3 ray_D [[direction]],
const float ray_tmin [[min_distance]],
const float ray_tmax [[max_distance]])
{
uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object);
const KernelCurveSegment segment = kernel_data_fetch(curve_segments, prim);
BoundingBoxIntersectionResult result;
result.accept = false;
result.continue_search = true;
result.distance = ray_tmax;
metalrt_intersection_curve(launch_params_metal,
payload,
object,
segment.prim,
segment.type,
ray_P,
ray_D,
# if defined(__METALRT_MOTION__)
payload.time,
# else
0.0f,
# endif
ray_tmin,
ray_tmax,
result);
return result;
}
[[intersection(bounding_box, triangle_data, METALRT_TAGS)]] BoundingBoxIntersectionResult
__intersection__curve_all_shadow(
constant KernelParamsMetal &launch_params_metal [[buffer(1)]],
ray_data MetalKernelContext::MetalRTIntersectionShadowPayload &payload [[payload]],
const uint object [[user_instance_id]],
const uint primitive_id [[primitive_id]],
const float3 ray_P [[origin]],
const float3 ray_D [[direction]],
const float ray_tmin [[min_distance]],
const float ray_tmax [[max_distance]])
{
uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object);
const KernelCurveSegment segment = kernel_data_fetch(curve_segments, prim);
BoundingBoxIntersectionResult result;
result.accept = false;
result.continue_search = true;
result.distance = ray_tmax;
metalrt_intersection_curve_shadow(launch_params_metal,
payload,
object,
segment.prim,
segment.type,
ray_P,
ray_D,
# if defined(__METALRT_MOTION__)
payload.time,
# else
0.0f,
# endif
ray_tmin,
ray_tmax,
result);
result.continue_search = metalrt_shadow_all_hit<METALRT_HIT_CURVE>(
launch_params_metal, payload, object, prim, float2(u, 0), ray_tmax, t, &ray);
result.accept = !result.continue_search;
return result;
}
@@ -685,13 +538,6 @@ ccl_device_inline void metalrt_intersection_point(
const float ray_tmax,
thread BoundingBoxIntersectionResult &result)
{
# ifdef __VISIBILITY_FLAG__
const uint visibility = payload.visibility;
if ((kernel_data_fetch(objects, object).visibility & visibility) == 0) {
return;
}
# endif
Intersection isect;
isect.t = ray_tmax;
@@ -703,10 +549,6 @@ ccl_device_inline void metalrt_intersection_point(
launch_params_metal, payload, object, prim, isect.u);
if (result.accept) {
result.distance = isect.t;
payload.u = isect.u;
payload.v = isect.v;
payload.prim = prim;
payload.type = type;
}
}
}
@@ -724,13 +566,6 @@ ccl_device_inline void metalrt_intersection_point_shadow(
const float ray_tmax,
thread BoundingBoxIntersectionResult &result)
{
# ifdef __VISIBILITY_FLAG__
const uint visibility = payload.visibility;
if ((kernel_data_fetch(objects, object).visibility & visibility) == 0) {
return;
}
# endif
Intersection isect;
isect.t = ray_tmax;
@@ -748,17 +583,21 @@ ccl_device_inline void metalrt_intersection_point_shadow(
}
}
[[intersection(bounding_box, triangle_data, METALRT_TAGS)]] BoundingBoxIntersectionResult
[[intersection(bounding_box, triangle_data, curve_data, METALRT_TAGS, extended_limits)]] BoundingBoxIntersectionResult
__intersection__point(constant KernelParamsMetal &launch_params_metal [[buffer(1)]],
ray_data MetalKernelContext::MetalRTIntersectionPayload &payload [[payload]],
const uint object [[user_instance_id]],
const uint object [[instance_id]],
const uint primitive_id [[primitive_id]],
const uint primitive_id_offset [[user_instance_id]],
const float3 ray_origin [[origin]],
const float3 ray_direction [[direction]],
# if defined(__METALRT_MOTION__)
const float time [[time]],
# endif
const float ray_tmin [[min_distance]],
const float ray_tmax [[max_distance]])
{
const uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object);
const uint prim = primitive_id + primitive_id_offset;
const int type = kernel_data_fetch(objects, object).primitive_type;
BoundingBoxIntersectionResult result;
@@ -774,7 +613,7 @@ __intersection__point(constant KernelParamsMetal &launch_params_metal [[buffer(1
ray_origin,
ray_direction,
# if defined(__METALRT_MOTION__)
payload.time,
time,
# else
0.0f,
# endif
@@ -785,18 +624,22 @@ __intersection__point(constant KernelParamsMetal &launch_params_metal [[buffer(1
return result;
}
[[intersection(bounding_box, triangle_data, METALRT_TAGS)]] BoundingBoxIntersectionResult
[[intersection(bounding_box, triangle_data, curve_data, METALRT_TAGS, extended_limits)]] BoundingBoxIntersectionResult
__intersection__point_shadow(constant KernelParamsMetal &launch_params_metal [[buffer(1)]],
ray_data MetalKernelContext::MetalRTIntersectionShadowPayload &payload
[[payload]],
const uint object [[user_instance_id]],
const uint object [[instance_id]],
const uint primitive_id [[primitive_id]],
const uint primitive_id_offset [[user_instance_id]],
const float3 ray_origin [[origin]],
const float3 ray_direction [[direction]],
# if defined(__METALRT_MOTION__)
const float time [[time]],
# endif
const float ray_tmin [[min_distance]],
const float ray_tmax [[max_distance]])
{
const uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object);
const uint prim = primitive_id + primitive_id_offset;
const int type = kernel_data_fetch(objects, object).primitive_type;
BoundingBoxIntersectionResult result;
@@ -812,7 +655,7 @@ __intersection__point_shadow(constant KernelParamsMetal &launch_params_metal [[b
ray_origin,
ray_direction,
# if defined(__METALRT_MOTION__)
payload.time,
time,
# else
0.0f,
# endif