Struct MPSRayIntersector

Source

pub struct MPSRayIntersector { /* private fields */ }

👎Deprecated

Available on crate features MPSCore and MPSKernel and MPSRayIntersector only.

Expand description

Performs intersection tests between rays and the geometry in an MPSAccelerationStructure

An MPSRayIntersector is used to schedule intersection tests between rays and geometry into an MTLCommandBuffer. First, create a raytracer with a Metal device. Then, configure the properties of the raytracer:

     id <MTLDevice> device = MTLCreateSystemDefaultDevice();
     id <MTLCommandQueue> commandQueue = [device newCommandQueue];
   
     MPSRayIntersector *raytracer = [[MPSRayIntersector alloc] initWithDevice:device];

     // Configure raytracer properties

Before scheduling intersection tests, an MPSAccelerationStructure must be created. The acceleration structure is built over geometry and is used to accelerate intersection testing. For example, to create a triangle acceleration structure, allocate an MPSTriangleAccelerationStructure object. Then, configure the properties of the acceleration structure. For example, triangle acceleration structures require a vertex buffer and a triangle count:

     MPSTriangleAccelerationStructure *accelerationStructure =
         [[MPSTriangleAccelerationStructure alloc] initWithDevice:device];
   
     accelerationStructure.vertexBuffer = vertexBuffer;
     accelerationStructure.triangleCount = triangleCount;

Acceleration structures must be built at least once before they are used for intersection testing, and must be rebuilt when the geometry changes. Rebuilding an acceleration structure is a time consuming operation, so an asynchronous version of this method is also available.

     [accelerationStructure rebuild];

The raytracer is then used to schedule intersection tests into an MTLCommandBuffer. Rays are provided in batches through a Metal buffer, and intersection results are returned through another Metal buffer in the same order, one intersection per ray.

There are several choices of ray data type controlled by the rayDataType property. The default ray data type is MPSRayOriginDirection, which includes just the ray origin direction. The other data types add support for minimum and maximum intersection distances and ray masks. These data types are available in the Metal Shading Language by including the MetalPerformanceShaders/MetalPerformanceShaders.h header. Additional application specific per-ray data can also be appended to the end of the ray data type using the rayStride property. This data will be ignored by the intersector.

If the rays were generated on the CPU:

     typedef MPSRayOriginDirection Ray;

     // Create a buffer to hold the rays
     id <MTLBuffer> rayBuffer = [device newBufferWithLength:sizeof(Ray) * rayCount options:0];

     // Copy the rays into the ray buffer
     memcpy(rayBuffer.contents, rays, sizeof(Ray) * rayCount);
   
     // Create a buffer to hold the intersections
     id <MTLBuffer> intersectionBuffer = [device newBufferWithLength:sizeof(Intersection) * rayCount
                                                             options:0];

It can be useful to prevent certain rays from participating in intersection testing. For example: rays which have bounced out of the scene in previous intersection tests. It may be more efficient to do this by compacting the ray buffer so that threads with invalid rays are not left idle during intersection testing. However, it can be more convenient to disable the ray in place. This can be done by setting most fields to invalid values. For example, setting the maximum distance to a negative value, setting the mask to zero, setting the direction to the zero vector, etc.

Finally, the intersection testing is encoded into an MTLCommandBuffer. There are two intersection types. The “nearest” intersection type returns the closest intersection along each ray. The “any” intersection type returns immediately when the first intersection is found. The “any” intersection type is useful for determining whether a point is visible from another point for, e.g., shadow rays or ambient occlusion rays and is typically much faster than the “nearest” intersection type.

     id <MTLCommandBuffer> commandBuffer = [commandQueue commandBuffer];
   
     [raytracer encodeIntersectionToCommandBuffer:commandBuffer
                                 intersectionType:MPSIntersectionTypeNearest
                                        rayBuffer:rayBuffer
                                  rayBufferOffset:0
                               intersectionBuffer:intersectionBuffer
                         intersectionBufferOffset:0
                                         rayCount:rayCount
                            accelerationStructure:accelerationStructure];

     [commandBuffer commit];

The intersection results are not available until the command buffer has finished executing on the GPU. It is not safe for the CPU to write or read the contents of the ray buffer, intersection buffer, vertex buffer, etc. until the command buffer has finished executing. Use the waitUntilCompleted or addCompletedHandler methods of the MTLCommandBuffer to block the CPU until the GPU has finished executing. Then retrieve the intersection results from the intersection buffer:

     typedef MPSIntersectionDistancePrimitiveIndexCoordinates Intersection;

     [commandBuffer waitUntilCompleted];

     Intersection *intersections = (Intersection *)intersectionBuffer.contents;

There are also several choices of intersection data type controlled by the intersectionDataType property. The default intersection data type is MPSIntersectionDistancePrimitiveIndexCoordinates, which includes the intersection distance, primitive index, and barycentric coordinates. The other data types remove the primitive index or barycentric coordinates, which can be used to reduce the memory and memory bandwidth usage of the intersection buffer. These data types are available in the Metal Shading Language by including the MetalPerformanceShaders/MetalPerformanceShaders.h header.

The intersection distance field is positive when an intersection has been found and negative when there is no intersection. When using the “nearest” intersection type, the intersection point is the ray origin plus the ray direction multiplied by the intersection distance. The other fields are not valid if there is no intersection. Only the intersection distance field is valid for the “any” intersection type, and the distance is either a negative or positive value to indicate an intersection or miss. It does not necessarily contain the actual intersection distance when using the “any” intersection type.

Asynchronous Raytracing: Copying rays and intersections to and from the CPU is expensive. Furthermore, generating rays and consuming intersections on the CPU causes the CPU and GPU to block each other. If the CPU must generate rays and consume intersections, it is better to add an asynchronous completion handler to the MTLCommandBuffer. The CPU can then proceed to do other useful work and will be notified when the GPU has finished executing. Use double or triple buffered ray and intersection buffers to avoid race conditions such as the CPU overwriting data the GPU may be reading. Then the CPU can safely write to one range of the buffer while the GPU reads from another range of the buffer. Once the GPU is done executing, the CPU and GPU can advance to the next range of the buffer. This method can be implemented using a completion handler and a semaphore:

     #define MAX_ASYNC_OPERATIONS 3

     // Initialization:

     // Create a semaphore with the maximum number of asynchronous operations in flight
     dispatch_semaphore_t asyncOperationSemaphore = dispatch_semaphore_create(MAX_ASYNC_OPERATIONS);

     // Create a ray and intersection buffer large enough for the maximum number of operations
     id <MTLBuffer> rayBuffer =
         [device newBufferWithLength:sizeof(Ray) * rayCount * MAX_ASYNC_OPERATIONS
                             options:0];

     id <MTLBuffer> intersectionBuffer =
         [device newBufferWithLength:sizeof(Intersection) * rayCount * MAX_ASYNC_OPERATIONS
                             options:0];

     NSUInteger asyncOperationIndex = 0;

     // Encode intersection testing:

     // Wait until there is a free buffer range
     dispatch_semaphore_wait(asyncOperationSemaphore, DISPATCH_TIME_FOREVER);

     // Copy rays into ray buffer
     NSUInteger rayBufferOffset = sizeof(Ray) * rayCount * asyncOperationIndex;
     NSUInteger intersectionBufferOffset = sizeof(Intersection) * rayCount * asyncOperationIndex;

     memcpy((uint8_t *)rayBuffer.contents + rayBufferOffset, rays, sizeof(Ray) * rayCount);

     // Advance the async operation index
     asyncOperationIndex = (asyncOperationIndex + 1) % MAX_ASYNC_OPERATIONS;

     // Create a command buffer
     id <MTLCommandBuffer> commandBuffer = [commandQueue commandBuffer];

     // Encode actual intersection work
     [raytracer encodeIntersectionToCommandBuffer:commandBuffer
                                 intersectionType:MPSIntersectionTypeNearest
                                        rayBuffer:rayBuffer
                                  rayBufferOffset:rayBufferOffset
                               intersectionBuffer:intersectionBuffer
                         intersectionBufferOffset:intersectionBufferOffset
                                         rayCount:rayCount
                            accelerationStructure:accelerationStructure];

     // Register a completion handler to run when the GPU finishes executing
     [commandBuffer addCompletedHandler:^(id <MTLCommandBuffer> commandBuffer) {
         Intersection *intersections = (Intersection *)((uint8_t *)intersectionBuffer.contents +
             intersectionBufferOffset);

         // Process intersections

         // Signal that the ray and intersection buffer ranges are now available for reuse
         dispatch_semaphore_signal(asyncOperationSemaphore);
     }];

     // Commit the command buffer to allow the GPU to start executing
     [commandBuffer commit];

GPU Driven Raytracing: Pipelining CPU and GPU work with asynchronous raytracing is better than allowing the CPU and GPU block each other, but it is even better to produce rays and consume intersections entirely on the GPU. This avoids the need to copy rays and intersections to and from the GPU and avoids any kind of CPU/GPU synchronization. To do this, encode compute kernels before and after intersection testing. By processing rays in parallel, the compute kernels may also be able to generate and consume rays faster than the CPU. The ray generation kernel typically produces rays according to some camera model, and the intersection consumption kernel typically updates the output buffer or texture according to some shading model.

Since the rays and intersections will never leave the GPU, store them in private Metal buffers that are allocated in GPU memory rather than system memory. Because the ray generation, intersection testing, and intersection consumption kernels are pipelined on the GPU, there is no need to double or triple buffer the ray or intersection buffers, which saves memory.

     id <MTLBuffer> rayBuffer =
         [device newBufferWithLength:sizeof(Ray) * rayCount
                             options:MTLResourceStorageModePrivate];
     id <MTLBuffer> intersectionBuffer =
         [device newBufferWithLength:sizeof(Intersection) * rayCount
                             options:MTLResourceStorageModePrivate];

     id <MTLCommandBuffer> commandBuffer = [commandQueue commandBuffer];

     // Generate rays
     id <MTLComputeCommandEncoder> encoder = [commandBuffer computeCommandEncoder];

     [encoder setBuffer:rayBuffer offset:0 atIndex:0];
     [encoder setBytes:&uniformData length:sizeof(uniformData) atIndex:1];

     [encoder setComputePipelineState:cameraPipeline];

     [encoder dispatchThreads:MTLSizeMake(rayCount, 1, 1)
        threadsPerThreadgroup:MTLSizeMake(64, 1, 1)];

     [encoder endEncoding];

     [raytracer encodeIntersectionToCommandBuffer:commandBuffer
                                 intersectionType:MPSIntersectionTypeNearest
                                        rayBuffer:rayBuffer
                                  rayBufferOffset:0
                               intersectionBuffer:intersectionBuffer
                         intersectionBufferOffset:0
                                         rayCount:rayCount
                            accelerationStructure:accelerationStructure];

     // Perform shading at intersections and update framebuffer texture
     encoder = [commandBuffer computeCommandEncoder];

     [encoder setBuffer:rayBuffer offset:0 atIndex:0];
     [encoder setBuffer:intersectionBuffer offset:0 atIndex:1];
     [encoder setBytes:&uniformData length:sizeof(uniformData) atIndex:2];

     [encoder setTexture:framebufferTexture atIndex:0];

     [encoder setComputePipelineState:shadingPipeline];

     [encoder dispatchThreads:MTLSizeMake(rayCount, 1, 1)
        threadsPerThreadgroup:MTLSizeMake(64, 1, 1)];

     [encoder endEncoding];

     [commandBuffer commit];

Note that the intersection consumption kernel can in turn produce new rays that can be passed back to the MPSRayIntersector. This technique can be used to implement iterative techniques such as progressive path tracing without leaving the GPU. For example, the shading kernel in the example above could produce both a secondary ray that will be passed back to the raytracer in the next iteration as well as a shadow ray that will be used to sample the direct lighting. A final kernel can consume the shadow ray intersections to accumulate lighting contributions into the framebuffer.

There is an alternative version of the intersection test encoding method that does not accept a literal ray count. The ray count is instead fetched indirectly by the GPU. For example, this can be combined with a parallel reduction on the GPU to compact the ray buffer after each iteration as rays bounce out of the scene or are absorbed. Alternatively, setting the maximum distance of a ray to a negative number indicates that the ray has become inactive and causes the raytracer to ignore the ray.

     [raytracer encodeIntersectionToCommandBuffer:commandBuffer
                                 intersectionType:MPSIntersectionTypeNearest
                                        rayBuffer:rayBuffer
                                  rayBufferOffset:0
                               intersectionBuffer:intersectionBuffer
                         intersectionBufferOffset:0
                                   rayCountBuffer:rayCountBuffer
                             rayCountBufferOffset:0
                            accelerationStructure:accelerationStructure];

Multi-GPU Raytracing: to implement multi-GPU raytracing, create the MPSRayIntersector and MPSAccelerationStructure objects first with one Metal device and copy them to the other Metal device(s). The raytracing process can then proceed independently on each GPU. For example, divide the output image into tiles or slices that are rendered independently. Then composite finished tiles or slices back together on one GPU and present the output image to the screen. The workload should be distributed across GPUs according to their performance to avoid a more powerful GPU idly waiting for a less powerful GPU to finish.

Acceleration Structure Serialization: MPSAccelerationStructure objects can be serialized and deserialized using the NSSecureCoding protocol. This can be used to build acceleration structures offline and reload them at runtime rather than building them from scratch.

Performance Guidelines:

For vertex buffers, ray buffers, intersection buffers, etc., use private or managed buffers rather than shared buffers when possible on discrete memory GPU architectures as they are much faster than fetching data over the PCIe bus. If the CPU only writes once to a ray buffer once and reads once from the intersection buffer, then a shared buffer may be acceptable and avoids extra copies to and from the GPU. However, it is generally preferable to generate and consume rays and intersections on the GPU instead, in which case a private buffer should be used. Vertex data is typically static and reused many times so it should be stored in private or managed buffers.
If the CPU must generate and consume rays and intersections, use double or triple buffering as described above. This avoids the CPU and GPU mutually blocking each other.
In general, disable any unused features such as ray masks, backface culling, etc. Enabling extra features increases the number of instructions and register usage of the ray intersection kernel(s), reducing intersection performance. For example, it may be more efficient to compute barycentric coordinates in your intersection consumption kernel rather getting them from the raytracer. Use of an index buffer may also reduce performance, so consider disabling the index buffer if there is enough memory available.
Try to submit rays in large batches. This amortizes the costs involved in dispatching work to the GPU and also allows the GPU to perform more effective latency hiding. Use the recommendedMinimumRayBatchSizeForRayCount method to get an estimate of the minimum recommended ray batch size. For this reason, small images or sample counts may not perform as well as large images or sample counts. Note, however, that submitting rays in very large batches can reduce the responsiveness of the system because the GPU will be busy for long periods. Experiment to find a balance between raytracing throughput and system responsiveness.
When possible, organize rays within a batch for spatial locality. Rays that originate at nearby points or are oriented in similar directions tend to access the same locations in memory and can therefore make more effective use of the GPU’s caches. For example, the camera rays associated with nearby pixels in the output image will likely originate at the same point and travel in very similar directions. Therefore, divide the output image into small tiles (e.g., 8x8). Rather than laying out all of the rays in the ray buffer in scanline order, first lay out the ray in scanline order within each tile, then lay out the tiles in scanline order or according to some space filling curve.
If CPU encode time is an issue, disable Metal API validation and enable MPSKernelOptionsSkipAPIValidation.
Choose the minimal ray and intersection data types for your use case. Loading and storing extra values such as ray masks or primitive indices can reduce raytracing performance, so use a simpler data type if they are not needed. For example, camera rays typically have no need for a maximum distance field, while shadow rays do.
Use MPSIntersectionTestTypeAny when possible: this is typically much faster than MPSIntersectionTestTypeNearest and can be used when you only need to check for binary visibility between two points such as shadow and ambient occlusion rays. Combine this with MPSRayDataTypeDistance to minimize memory bandwidth usage.
Try to keep the geometry, textures, ray buffers, etc. within the Metal device’s recommended working set size. Paging data into GPU memory can significantly reduce raytracing performance.
Changes to MPSRayIntersector properties can trigger internal pipeline compilations when intersection tests are next encoded. If you need to avoid hitches due to pipeline compilation, encode a small ray intersection with each raytracer configuration you will use at encode-time. This creates and caches the corresponding pipelines.
Disable rays which should not participate in intersection testing. This can be done either by compacting the ray buffer such that it only contains valid rays, or by setting fields of the ray struct to invalid values. For example, setting the maximum distance to a negative value, setting the mask to zero, setting the direction to the zero vector, etc. In particular, rays should NOT be disabled using schemes such as moving their origin outside the scene. These rays will still partially traverse the acceleration structure, potentially evicting data from the cache which could have been used by valid rays. Note that it is preferable to provide only valid rays so that threads are not left idle if their rays are found to be invalid, but it can be convenient to disable rays in place in the ray buffer.

See MPSAccelerationStructure and MPSInstanceAccelerationStructure for more performance guidelines.

Thread Safety: MPSRayIntersectors are generally not thread safe: changing properties and encoding intersection tests from multiple threads result in undefined behavior. Instead, multiple threads should copy or create their own MPSRayIntersectors.

Struct MPSRayIntersector Copy item path

Implementations§

impl MPSRayIntersector

pub unsafe fn cullMode(&self) -> MTLCullMode

pub unsafe fn setCullMode(&self, cull_mode: MTLCullMode)

pub unsafe fn frontFacingWinding(&self) -> MTLWinding

pub unsafe fn setFrontFacingWinding(&self, front_facing_winding: MTLWinding)

pub unsafe fn triangleIntersectionTestType( &self, ) -> MPSTriangleIntersectionTestType

pub unsafe fn setTriangleIntersectionTestType( &self, triangle_intersection_test_type: MPSTriangleIntersectionTestType, )

pub unsafe fn boundingBoxIntersectionTestType( &self, ) -> MPSBoundingBoxIntersectionTestType

pub unsafe fn setBoundingBoxIntersectionTestType( &self, bounding_box_intersection_test_type: MPSBoundingBoxIntersectionTestType, )

pub unsafe fn rayMaskOptions(&self) -> MPSRayMaskOptions

pub unsafe fn setRayMaskOptions(&self, ray_mask_options: MPSRayMaskOptions)

pub unsafe fn rayMaskOperator(&self) -> MPSRayMaskOperator

pub unsafe fn setRayMaskOperator(&self, ray_mask_operator: MPSRayMaskOperator)

pub unsafe fn rayStride(&self) -> NSUInteger

pub unsafe fn setRayStride(&self, ray_stride: NSUInteger)

pub unsafe fn intersectionStride(&self) -> NSUInteger

pub unsafe fn setIntersectionStride(&self, intersection_stride: NSUInteger)

pub unsafe fn rayDataType(&self) -> MPSRayDataType

pub unsafe fn setRayDataType(&self, ray_data_type: MPSRayDataType)

pub unsafe fn intersectionDataType(&self) -> MPSIntersectionDataType

pub unsafe fn setIntersectionDataType( &self, intersection_data_type: MPSIntersectionDataType, )

pub unsafe fn rayIndexDataType(&self) -> MPSDataType

pub unsafe fn setRayIndexDataType(&self, ray_index_data_type: MPSDataType)

pub unsafe fn rayMask(&self) -> c_uint

pub unsafe fn setRayMask(&self, ray_mask: c_uint)

pub unsafe fn init(this: Allocated<Self>) -> Retained<Self>

pub unsafe fn initWithDevice( this: Allocated<Self>, device: &ProtocolObject<dyn MTLDevice>, ) -> Retained<Self>

pub unsafe fn initWithCoder_device( this: Allocated<Self>, a_decoder: &NSCoder, device: &ProtocolObject<dyn MTLDevice>, ) -> Option<Retained<Self>>

§Safety

pub unsafe fn copyWithZone_device( &self, zone: *mut NSZone, device: Option<&ProtocolObject<dyn MTLDevice>>, ) -> Retained<Self>

§Safety

pub unsafe fn recommendedMinimumRayBatchSizeForRayCount( &self, ray_count: NSUInteger, ) -> NSUInteger

pub unsafe fn encodeWithCoder(&self, coder: &NSCoder)

§Safety

§Safety

§Safety

§Safety

§Safety

§Safety

impl MPSRayIntersector

pub unsafe fn initWithCoder( this: Allocated<Self>, a_decoder: &NSCoder, ) -> Option<Retained<Self>>

§Safety

impl MPSRayIntersector

pub unsafe fn new() -> Retained<Self>

Methods from Deref<Target = MPSKernel>§

pub unsafe fn options(&self) -> MPSKernelOptions

pub unsafe fn setOptions(&self, options: MPSKernelOptions)

pub unsafe fn device(&self) -> Retained<ProtocolObject<dyn MTLDevice>>

pub unsafe fn label(&self) -> Option<Retained<NSString>>

pub unsafe fn setLabel(&self, label: Option<&NSString>)

pub unsafe fn copyWithZone_device( &self, zone: *mut NSZone, device: Option<&ProtocolObject<dyn MTLDevice>>, ) -> Retained<Self>

§Safety

Methods from Deref<Target = NSObject>§

pub fn doesNotRecognizeSelector(&self, sel: Sel) -> !

Methods from Deref<Target = AnyObject>§

pub fn class(&self) -> &'static AnyClass

§Panics

§Example

pub unsafe fn get_ivar<T>(&self, name: &str) -> &Twhere T: Encode,

§Safety

pub fn downcast_ref<T>(&self) -> Option<&T>where T: DowncastTarget,

§Mutable classes

§Generic classes

§Panics

§Examples

Trait Implementations§

impl AsRef<AnyObject> for MPSRayIntersector

fn as_ref(&self) -> &AnyObject

impl AsRef<MPSKernel> for MPSRayIntersector

fn as_ref(&self) -> &MPSKernel

impl AsRef<MPSRayIntersector> for MPSRayIntersector

fn as_ref(&self) -> &Self

impl AsRef<NSObject> for MPSRayIntersector

fn as_ref(&self) -> &NSObject

impl Borrow<AnyObject> for MPSRayIntersector

fn borrow(&self) -> &AnyObject

impl Borrow<MPSKernel> for MPSRayIntersector

fn borrow(&self) -> &MPSKernel

Struct MPSRayIntersector

pub unsafe fn get_ivar<T>(&self, name: &str) -> &T
where T: Encode,

pub fn downcast_ref<T>(&self) -> Option<&T>
where T: DowncastTarget,

fn hash_slice<H>(data: &[Self], state: &mut H)
where H: Hasher, Self: Sized,

fn retain(&self) -> Retained<Self>
where Self: Sized,

unsafe fn encodeWithCoder(&self, coder: &NSCoder)
where Self: Sized + Message,

unsafe fn initWithCoder( this: Allocated<Self>, coder: &NSCoder, ) -> Option<Retained<Self>>
where Self: Sized + Message,

fn copy(&self) -> Retained<Self::Result>
where Self: Sized + Message + CopyingHelper,

unsafe fn copyWithZone(&self, zone: *mut NSZone) -> Retained<Self::Result>
where Self: Sized + Message + CopyingHelper,

fn isEqual(&self, other: Option<&AnyObject>) -> bool
where Self: Sized + Message,

fn hash(&self) -> usize
where Self: Sized + Message,

fn isKindOfClass(&self, cls: &AnyClass) -> bool
where Self: Sized + Message,

fn is_kind_of<T>(&self) -> bool
where T: ClassType, Self: Sized + Message,

fn isMemberOfClass(&self, cls: &AnyClass) -> bool
where Self: Sized + Message,

fn respondsToSelector(&self, aSelector: Sel) -> bool
where Self: Sized + Message,

fn conformsToProtocol(&self, aProtocol: &AnyProtocol) -> bool
where Self: Sized + Message,

fn description(&self) -> Retained<NSObject>
where Self: Sized + Message,

fn debugDescription(&self) -> Retained<NSObject>
where Self: Sized + Message,

fn isProxy(&self) -> bool
where Self: Sized + Message,

fn retainCount(&self) -> usize
where Self: Sized + Message,

fn supportsSecureCoding() -> bool
where Self: Sized + ClassType,