Skip to content

Comments

Fix #10092: empty ray payload struct generates invalid CUDA code#10100

Open
szihs wants to merge 1 commit intoshader-slang:masterfrom
szihs:dev/haaggarwal/fix-empty-raypayload
Open

Fix #10092: empty ray payload struct generates invalid CUDA code#10100
szihs wants to merge 1 commit intoshader-slang:masterfrom
szihs:dev/haaggarwal/fix-empty-raypayload

Conversation

@szihs
Copy link
Collaborator

@szihs szihs commented Feb 20, 2026

Add non-template optixTrace/optixTraverse overloads for the empty payload case. When Slang's type legalization eliminates an empty payload struct, the generated CUDA code omits the payload argument, which failed to match any existing prelude wrapper.

…DA code

Add non-template optixTrace/optixTraverse overloads for the empty payload
case. When Slang's type legalization eliminates an empty payload struct,
the generated CUDA code omits the payload argument, which failed to match
any existing prelude wrapper.
Copilot AI review requested due to automatic review settings February 20, 2026 09:56
@szihs szihs requested a review from a team as a code owner February 20, 2026 09:56
@szihs szihs requested review from bmillsNV and removed request for a team February 20, 2026 09:56
@coderabbitai
Copy link

coderabbitai bot commented Feb 20, 2026

No actionable comments were generated in the recent review. 🎉


📝 Walkthrough

Walkthrough

Added non-template overloads for OptiX tracing, traversal, and invoke functions in the CUDA prelude header to support empty payload scenarios. These wrappers enable code generated when payloads are eliminated by type legalization to compile without providing payload arguments.

Changes

Cohort / File(s) Summary
OptiX Empty Payload Overloads
prelude/slang-cuda-prelude.h
Added four non-template function overloads: two variants of optixTrace, one optixTraverse, and one optixInvoke, each handling cases where payloads are absent by forwarding calls to existing functions with appropriate parameter omissions.

Estimated code review effort

🎯 1 (Trivial) | ⏱️ ~3 minutes

Poem

🐰 Empty payloads, they said, won't compile,
So we crafted some wrappers with OptiX style,
No data to pass? No problem at all!
Our overloads answer the payload-less call! 🌟

🚥 Pre-merge checks | ✅ 2 | ❌ 1

❌ Failed checks (1 warning)

Check name Status Explanation Resolution
Docstring Coverage ⚠️ Warning Docstring coverage is 0.00% which is insufficient. The required threshold is 80.00%. Write docstrings for the functions missing them to satisfy the coverage threshold.
✅ Passed checks (2 passed)
Check name Status Explanation
Title check ✅ Passed The title clearly and specifically describes the main change: fixing CUDA code generation for empty ray payload structs by adding non-template overloads.
Description check ✅ Passed The description is directly related to the changeset, explaining the problem (empty payload elimination) and the solution (non-template overloads for matching generated calls).

✏️ Tip: You can configure your own custom pre-merge checks in the settings.

✨ Finishing Touches
  • 📝 Generate docstrings (stacked PR)
  • 📝 Generate docstrings (commit on current branch)
🧪 Generate unit tests (beta)
  • Create PR with unit tests
  • Post copyable unit tests in a comment

Comment @coderabbitai help to get the list of available commands and usage tips.

Copy link
Contributor

Copilot AI left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Pull request overview

Fixes CUDA/OptiX codegen for the case where Slang’s legalization removes an empty ray payload struct, causing generated CUDA to call optixTrace / optixTraverse without a payload argument that previously didn’t match any Slang-provided wrapper overload.

Changes:

  • Added a non-template optixTrace(...) overload that accepts no payload argument and forwards to the native OptiX intrinsic.
  • Added a non-template optixTraverse(...) overload that accepts no payload argument (but keeps the hitObj parameter for signature compatibility) and forwards to the native OptiX intrinsic.

Comment on lines +4656 to +4668
// Non-template overload for empty payload case.
// When Slang's type legalization eliminates an empty payload struct,
// the generated code calls optixTrace without a payload argument.
__forceinline__ __device__ void optixTrace(
OptixTraversableHandle AccelerationStructure,
uint32_t RayFlags,
uint32_t InstanceInclusionMask,
uint32_t RayContributionToHitGroupIndex,
uint32_t MultiplierForGeometryContributionToHitGroupIndex,
uint32_t MissShaderIndex,
RayDesc Ray)
{
optixTrace(
Copy link

Copilot AI Feb 20, 2026

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Please add a regression test that exercises the empty ray payload case on the CUDA target (e.g., a [raypayload] struct with no fields used with TraceRay/HitObject::TraceRay) and filecheck that the generated CUDA calls the new no-payload optixTrace/optixTraverse overload(s). This change fixes a codegen mismatch, but without a test it’s easy for future prelude/codegen changes to reintroduce the invalid CUDA call signature from #10092.

Copilot generated this review using guidance from repository custom instructions.
@szihs szihs added the pr: non-breaking PRs without breaking changes label Feb 20, 2026
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

pr: non-breaking PRs without breaking changes

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant