diff --git a/intern/cycles/blender/addon/properties.py b/intern/cycles/blender/addon/properties.py index 9177eeaef1b..75b44215536 100644 --- a/intern/cycles/blender/addon/properties.py +++ b/intern/cycles/blender/addon/properties.py @@ -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() diff --git a/intern/cycles/blender/python.cpp b/intern/cycles/blender/python.cpp index 7e6ed085551..42a26e2f006 100644 --- a/intern/cycles/blender/python.cpp +++ b/intern/cycles/blender/python.cpp @@ -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); } diff --git a/intern/cycles/device/metal/bvh.h b/intern/cycles/device/metal/bvh.h index 4882aa569d4..ce3f770c0e7 100644 --- a/intern/cycles/device/metal/bvh.h +++ b/intern/cycles/device/metal/bvh.h @@ -22,7 +22,9 @@ class BVHMetal : public BVH { API_AVAILABLE(macos(11.0)) vector> blas_array; - vector blas_lookup; + + API_AVAILABLE(macos(11.0)) + vector> unique_blas_array; bool motion_blur = false; diff --git a/intern/cycles/device/metal/bvh.mm b/intern/cycles/device/metal/bvh.mm index 267d8ab9ee5..c1b838e35ed 100644 --- a/intern/cycles/device/metal/bvh.mm +++ b/intern/cycles/device/metal/bvh.mm @@ -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(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 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 cpBuffer = nil; + id radiusBuffer = nil; + id idxBuffer = nil; MTLAccelerationStructureGeometryDescriptor *geomDesc; if (motion_blur) { - std::vector 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 &radiuses = hair->get_curve_radius(); + + /* Gather the curve geometry. */ + std::vector cpData; + std::vector idxData; + std::vector radiusData; + cpData.reserve(numKeys); + radiusData.reserve(numKeys); + + std::vector 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 cp_ptrs; + std::vector 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 &radiuses = hair->get_curve_radius(); + + /* Gather the curve geometry. */ + std::vector cpData; + std::vector idxData; + std::vector 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 /*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(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(const_cast(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(const_cast(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( + const_cast(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 blas in all_blas) { - blas_array.push_back(blas); - } + + unique_blas_array.clear(); + unique_blas_array.reserve(all_blas.count); + [all_blas enumerateObjectsUsingBlock:^( + id blas, NSUInteger, BOOL *) { + unique_blas_array.push_back(blas); + }]; return true; } diff --git a/intern/cycles/device/metal/device.mm b/intern/cycles/device/metal/device.mm index 188e90a8fc5..f3e951a0b26 100644 --- a/intern/cycles/device/metal/device.mm +++ b/intern/cycles/device/metal/device.mm @@ -62,12 +62,17 @@ void device_metal_info(vector &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++; diff --git a/intern/cycles/device/metal/device_impl.h b/intern/cycles/device/metal/device_impl.h index a83eb79422b..5dcc9d545c1 100644 --- a/intern/cycles/device/metal/device_impl.h +++ b/intern/cycles/device/metal/device_impl.h @@ -82,7 +82,6 @@ class MetalDevice : public Device { /* BLAS encoding & lookup */ id mtlBlasArgEncoder = nil; id blas_buffer = nil; - id blas_lookup_buffer = nil; bool use_metalrt = false; MetalPipelineType kernel_specialization_level = PSO_GENERIC; diff --git a/intern/cycles/device/metal/device_impl.mm b/intern/cycles/device/metal/device_impl.mm index 657a3da25e9..1a7ce6e0ddc 100644 --- a/intern/cycles/device/metal/device_impl.mm +++ b/intern/cycles/device/metal/device_impl.mm @@ -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)]; } } } diff --git a/intern/cycles/device/metal/kernel.h b/intern/cycles/device/metal/kernel.h index 36b5669c5a8..82dbb878089 100644 --- a/intern/cycles/device/metal/kernel.h +++ b/intern/cycles/device/metal/kernel.h @@ -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 diff --git a/intern/cycles/device/metal/kernel.mm b/intern/cycles/device/metal/kernel.mm index d1b7f96a464..c8b2ba4379a 100644 --- a/intern/cycles/device/metal/kernel.mm +++ b/intern/cycles/device/metal/kernel.mm @@ -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 point_intersect_default = nil; id 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]]; diff --git a/intern/cycles/device/metal/queue.mm b/intern/cycles/device/metal/queue.mm index 71c0eb7cdf1..0d252ce33e4 100644 --- a/intern/cycles/device/metal/queue.mm +++ b/intern/cycles/device/metal/queue.mm @@ -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]; } } diff --git a/intern/cycles/kernel/device/metal/bvh.h b/intern/cycles/kernel/device/metal/bvh.h index 658b3244398..870ccfa6968 100644 --- a/intern/cycles/kernel/device/metal/bvh.h +++ b/intern/cycles/kernel/device/metal/bvh.h @@ -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 diff --git a/intern/cycles/kernel/device/metal/compat.h b/intern/cycles/kernel/device/metal/compat.h index fcd50ced5f1..47b61af40c7 100644 --- a/intern/cycles/kernel/device/metal/compat.h +++ b/intern/cycles/kernel/device/metal/compat.h @@ -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_as_type; -typedef intersection_function_table metalrt_ift_type; -typedef metal::raytracing::intersector metalrt_intersector_type; +typedef intersection_function_table + metalrt_ift_type; +typedef metal::raytracing::intersector + metalrt_intersector_type; # if defined(__METALRT_MOTION__) typedef acceleration_structure metalrt_blas_as_type; -typedef intersection_function_table metalrt_blas_ift_type; -typedef metal::raytracing::intersector - metalrt_blas_intersector_type; +typedef intersection_function_table + metalrt_blas_ift_type; +typedef metal::raytracing:: + intersector + metalrt_blas_intersector_type; # else typedef acceleration_structure<> metalrt_blas_as_type; -typedef intersection_function_table metalrt_blas_ift_type; -typedef metal::raytracing::intersector metalrt_blas_intersector_type; +typedef intersection_function_table + metalrt_blas_ift_type; +typedef metal::raytracing::intersector + 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 }; diff --git a/intern/cycles/kernel/device/metal/kernel.metal b/intern/cycles/kernel/device/metal/kernel.metal index 91b34b38da2..42f80431f29 100644 --- a/intern/cycles/kernel/device/metal/kernel.metal +++ b/intern/cycles/kernel/device/metal/kernel.metal @@ -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 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( - launch_params_metal, payload, payload.local_object, primitive_id, barycentrics, ray_tmax); + return metalrt_local_hit( + 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( - launch_params_metal, payload, instance_id, primitive_id, barycentrics, ray_tmax); + return metalrt_local_hit( + 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( 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( + uint prim = primitive_id + primitive_id_offset; + PrimitiveIntersectionResult result = + metalrt_visibility_test( 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( - 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( - 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( + 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( + 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