Slang OptiX anyhit loses payload compared to cuda equivalent when using IgnoreHit

[oliver-batchelor]

I wrote a slang version of a cuda example which collects all the intersections in a hit buffer. The main difference seems to be in how the payload is handled, in cuda a pointer to the local stack is used, then the anyhit shader will write directly to that.

In slang there's an inout parameter which is mutated. However, every time the anyhit is called the payload seems to be reset to it's original state.

Some investigation - if IgnoreHit is not called, then ONE hit is registered, but when IgnoreHit is called, any mutations to the payload seem to disappear.

#define INF 1e10

#ifndef MAX_HITS
  #define MAX_HITS 16
#endif


struct RayHit {
    float t;
    int id;
};

struct CameraParams {
    float2 principle_point;
    float2 focal_length;

    float3x3 rotation; // 3x3 matrix world_t_camera
    float3 origin;

    int2 image_size;
    float near, far; // near and far clipping planes

    float3 direction(float2 pixel) {
        var uv = float3((pixel - principle_point) / focal_length, 1.0f);
        return normalize(mul(uv, rotation));
    };
};

CameraParams camera;
RWStructuredBuffer<RayHit> output_hits;
uniform RaytracingAccelerationStructure sceneBVH;

[payload]
struct HitBuffer<let N : int = MAX_HITS> {
    RayHit buffer[N];

    // Function to initialize the HitBuffer with default values
    [mutating]
    void initialize() {
        for (int i = 0; i < N; i++) {
            this.buffer[i] = RayHit( { INF, -1 }); // Set a large t value (e.g., 1e10) and id = -1
        }
    }

    RayHit hit(int i)  {
        return buffer[i];
    }

    float last() {
        return buffer[N - 1].t;
    }

    [mutating]
    int insert(RayHit hit) {
      int index = 0;

      var current = hit;  
      for (int i = 0; i < N; ++i) {
          if (current.t < buffer[i].t) {
          index = min(index, i);

          let temp = buffer[i];
          buffer[i] = current;
          current = temp;
        }
      }

      return index;
    }
};


[shader("raygen")]
void rayGenShader()
{
    uint2 p = DispatchRaysIndex().xy;
    if (p.x >= (int)camera.image_size.x || p.y >= (int)camera.image_size.y) return;

    var pixel = float2(p.x + 0.5f, p.y + 0.5f);
    var dir = camera.direction(pixel);

    RayDesc ray;
    ray.Origin = camera.origin;
    ray.Direction = dir;
    ray.TMin = camera.near;
    ray.TMax = camera.far;

    var hit_buffer : HitBuffer<MAX_HITS>;
    hit_buffer.initialize();

    TraceRay<HitBuffer<MAX_HITS>>(sceneBVH, RAY_FLAG_CULL_BACK_FACING_TRIANGLES, 
      ~0, 0, 0, 0, ray, hit_buffer);

    int offset = (p.y * camera.image_size.x + p.x) * MAX_HITS;
    for (int i = 0; i < MAX_HITS; i++) {
        output_hits[offset + i] = hit_buffer.hit(i);
    }
    
}

[shader("anyhit")]
void anyHitShader(inout HitBuffer<MAX_HITS> buffer)
{
    RayHit hit = { RayTCurrent(), InstanceID()};
    int index = buffer.insert(hit);

    if (hit.t < buffer.last()) {
    IgnoreHit();
    }
}

[shader("miss")]
void missShader(inout HitBuffer<MAX_HITS> buffer) {
}

[oliver-batchelor]

Moreover, if I compile it to cuda and examine the output, I can see it uses the same trick to unpack a local pointer, but it does this TWICE, the first time (line 102) making a copy on the stack,

Then, when everything is done (line 107) it takes the copy and writes it back over the mutations performed in the function.

extern "C" __global__ void __anyhit__anyHitShader()
{

#line 102
    HitBuffer_0 * _S11 = (HitBuffer_0 *)getOptiXRayPayloadPtr();

#line 102
    HitBuffer_0 _S12 = *_S11;

#line 102
    HitBuffer_0 * _S13 = (HitBuffer_0 *)getOptiXRayPayloadPtr();

    float _S14 = optixGetRayTmax();

#line 104
    uint _S15 = optixGetInstanceId();

#line 104
    RayHit_0 hit_1 = { _S14, int(_S15) };
    int _S16 = HitBuffer_insert_0(&_S12, hit_1);

    if(_S14 < HitBuffer_last_0(_S12))
    {

#line 108
        optixIgnoreIntersection();

#line 107
    }

#line 107
    *_S13 = _S12;


    return;
}

[oliver-batchelor]

Some investigation reveals the problem is that IgnoreHit() does not return and therefore the last line
*_S13 = _S12; is never called. So it's somewhat like an "exception".

The state is copied in S12 which is operated on, then copied back to S13 at the end on line 107, so mutations of S12 don't change the state directly.

I will attempt to work on a minimal reproducer