Support for Payload Access Qualifiers (#3448) (#6595)

* Add support for Ray Payload Access Qualifiers (PAQs) (#3448)

- Added [raypayload] attribute for struct declarations
- Implemented field validation requiring read/write access qualifiers
- Added diagnostic error for missing qualifiers
- Enabled PAQs in DXC compiler and HLSL emission
- Added new test demonstrating PAQ syntax
- Implemented proper handling of ray payload attributes in IR generation

* format code

* Cleanup: Remove unused vars

* Add check to enablePAQ only for profile >= lib_6_7

* Review Fix - Add PAQ support for DX Raytracing

add enablePAQ flag to DownstreamCompileOpitons, improve PAQ handling
update raypayload-attribute-paq.slang to ensure hlsl and dxil is
validated

* Add diagnostic test for missing paq for lib_6_7

Compile using `-disable-payload-qualifiers` aka lib_6_6 profile
raypayload-attribute-no-struct.slang and
raypayload-attribute.slang

---------

Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
Co-authored-by: Ellie Hermaszewska <ellieh@nvidia.com>

18 files changed, 196 insertions, 14 deletions

diff --git a/source/compiler-core/slang-downstream-compiler.h b/source/compiler-core/slang-downstream-compiler.h
index 5365b9839..c23a6eff0 100644
--- a/source/compiler-core/slang-downstream-compiler.h
+++ b/source/compiler-core/slang-downstream-compiler.h
@@ -260,6 +260,9 @@ struct DownstreamCompileOptions
     /// Profile name to use, only required for compiles that need to compile against a a specific
     /// profiles. Profile names are tied to compilers and targets.
     TerminatedCharSlice profileName;
+    // According to DirectX Raytracing Specification, PAQs are supported in Shader Model 6.7 and
+    // above
+    bool enablePAQ = false;
 
     /// The stage being compiled for
     SlangStage stage = SLANG_STAGE_NONE;
diff --git a/source/compiler-core/slang-dxc-compiler.cpp b/source/compiler-core/slang-dxc-compiler.cpp
index 065ee4145..0d4bc0a59 100644
--- a/source/compiler-core/slang-dxc-compiler.cpp
+++ b/source/compiler-core/slang-dxc-compiler.cpp
@@ -479,8 +479,7 @@ SlangResult DXCDownstreamCompiler::compile(const CompileOptions& inOptions, IArt
         args.add(compilerSpecific[i]);
     }
 
-    // This can be re-enabled when we add PAQs: https://github.com/shader-slang/slang/issues/3448
-    const bool enablePAQs = false;
+    bool enablePAQs = options.enablePAQ;
     if (!enablePAQs)
         args.add(L"-disable-payload-qualifiers");
     else
diff --git a/source/slang/core.meta.slang b/source/slang/core.meta.slang
index e2fb8bbf2..481aba191 100644
--- a/source/slang/core.meta.slang
+++ b/source/slang/core.meta.slang
@@ -4214,3 +4214,6 @@ attribute_syntax [RequireFullQuads] : RequireFullQuadsAttribute;
 __generic<T>
 typealias NodePayloadPtr = Ptr<T, $( (uint64_t)AddressSpace::NodePayloadAMDX)>;
 
+__attributeTarget(StructDecl)
+attribute_syntax [raypayload] : RayPayloadAttribute;
+
diff --git a/source/slang/slang-ast-modifier.h b/source/slang/slang-ast-modifier.h
index 5f9ccb5bb..86c1b556c 100644
--- a/source/slang/slang-ast-modifier.h
+++ b/source/slang/slang-ast-modifier.h
@@ -1698,6 +1698,16 @@ class PayloadAttribute : public Attribute
     SLANG_AST_CLASS(PayloadAttribute)
 };
 
+/// A `[raypayload]` attribute indicates that a `struct` type will be used as
+/// a ray payload for `TraceRay()` calls, and thus also as input/output
+/// for shaders in the ray tracing pipeline that might be invoked for
+/// such a ray.
+///
+class RayPayloadAttribute : public Attribute
+{
+    SLANG_AST_CLASS(RayPayloadAttribute)
+};
+
 /// A `[deprecated("message")]` attribute indicates the target is
 /// deprecated.
 /// A compiler warning including the message will be raised if the
diff --git a/source/slang/slang-check-decl.cpp b/source/slang/slang-check-decl.cpp
index 4ab909118..21a16cae5 100644
--- a/source/slang/slang-check-decl.cpp
+++ b/source/slang/slang-check-decl.cpp
@@ -12496,6 +12496,12 @@ void SemanticsDeclAttributesVisitor::visitStructDecl(StructDecl* structDecl)
         }
     }
 
+    // Check if this is a ray payload struct and validate field access qualifiers
+    if (structDecl->findModifier<RayPayloadAttribute>())
+    {
+        checkRayPayloadStructFields(structDecl);
+    }
+
     int backingWidth = 0;
     [[maybe_unused]] int totalWidth = 0;
     struct BitFieldInfo
diff --git a/source/slang/slang-check-impl.h b/source/slang/slang-check-impl.h
index f7681ba45..8a1e79ce8 100644
--- a/source/slang/slang-check-impl.h
+++ b/source/slang/slang-check-impl.h
@@ -2835,6 +2835,8 @@ public:
     bool isCStyleType(Type* type, HashSet<Type*>& isVisit);
 
     void addVisibilityModifier(Decl* decl, DeclVisibility vis);
+
+    void checkRayPayloadStructFields(StructDecl* structDecl);
 };
 
 
diff --git a/source/slang/slang-check-modifier.cpp b/source/slang/slang-check-modifier.cpp
index 741823a65..d94c77d6a 100644
--- a/source/slang/slang-check-modifier.cpp
+++ b/source/slang/slang-check-modifier.cpp
@@ -1413,9 +1413,20 @@ bool isModifierAllowedOnDecl(bool isGLSLInput, ASTNodeType modifierType, Decl* d
     case ASTNodeType::ConstRefModifier:
     case ASTNodeType::GLSLBufferModifier:
     case ASTNodeType::GLSLPatchModifier:
+        return (as<VarDeclBase>(decl) && isGlobalDecl(decl)) || as<ParamDecl>(decl) ||
+               as<GLSLInterfaceBlockDecl>(decl);
     case ASTNodeType::RayPayloadAccessSemantic:
     case ASTNodeType::RayPayloadReadSemantic:
     case ASTNodeType::RayPayloadWriteSemantic:
+        // Allow on struct fields if the parent struct has the [raypayload] attribute
+        if (auto varDecl = as<VarDeclBase>(decl))
+        {
+            if (auto structDecl = as<StructDecl>(varDecl->parentDecl))
+            {
+                if (structDecl->findModifier<RayPayloadAttribute>())
+                    return true;
+            }
+        }
         return (as<VarDeclBase>(decl) && isGlobalDecl(decl)) || as<ParamDecl>(decl) ||
                as<GLSLInterfaceBlockDecl>(decl);
 
@@ -2179,5 +2190,36 @@ void SemanticsVisitor::checkModifiers(ModifiableSyntaxNode* syntaxNode)
     postProcessingOnModifiers(syntaxNode->modifiers);
 }
 
+void SemanticsVisitor::checkRayPayloadStructFields(StructDecl* structDecl)
+{
+    // Only check structs with the [raypayload] attribute
+    if (!structDecl->findModifier<RayPayloadAttribute>())
+    {
+        return;
+    }
+
+    // Check each field in the struct
+    for (auto member : structDecl->members)
+    {
+        auto fieldVarDecl = as<VarDeclBase>(member);
+        if (!fieldVarDecl)
+        {
+            continue;
+        }
+
+        bool hasReadModifier = fieldVarDecl->findModifier<RayPayloadReadSemantic>() != nullptr;
+        bool hasWriteModifier = fieldVarDecl->findModifier<RayPayloadWriteSemantic>() != nullptr;
+
+        if (!hasReadModifier && !hasWriteModifier)
+        {
+            // Emit the diagnostic error
+            getSink()->diagnose(
+                fieldVarDecl,
+                Diagnostics::rayPayloadFieldMissingAccessQualifiers,
+                fieldVarDecl->getName());
+        }
+    }
+}
+
 
 } // namespace Slang
diff --git a/source/slang/slang-compiler.cpp b/source/slang/slang-compiler.cpp
index 3839e0722..55f3846af 100644
--- a/source/slang/slang-compiler.cpp
+++ b/source/slang/slang-compiler.cpp
@@ -1732,6 +1732,11 @@ SlangResult CodeGenContext::emitWithDownstreamForEntryPoints(ComPtr<IArtifact>&
     options.libraries = SliceUtil::asSlice(libraries);
     options.libraryPaths = allocator.allocate(libraryPaths);
 
+    if (m_targetProfile.getFamily() == ProfileFamily::DX)
+    {
+        options.enablePAQ = m_targetProfile.getVersion() >= ProfileVersion::DX_6_7;
+    }
+
     // Compile
     ComPtr<IArtifact> artifact;
     auto downstreamStartTime = std::chrono::high_resolution_clock::now();
diff --git a/source/slang/slang-compiler.h b/source/slang/slang-compiler.h
index 18192678a..8a9b8985a 100644
--- a/source/slang/slang-compiler.h
+++ b/source/slang/slang-compiler.h
@@ -2812,7 +2812,9 @@ public:
     };
 
     CodeGenContext(Shared* shared)
-        : m_shared(shared), m_targetFormat(shared->targetProgram->getTargetReq()->getTarget())
+        : m_shared(shared)
+        , m_targetFormat(shared->targetProgram->getTargetReq()->getTarget())
+        , m_targetProfile(shared->targetProgram->getOptionSet().getProfile())
     {
     }
 
@@ -2909,6 +2911,7 @@ public:
 
 protected:
     CodeGenTarget m_targetFormat = CodeGenTarget::Unknown;
+    Profile m_targetProfile;
     ExtensionTracker* m_extensionTracker = nullptr;
 
     /// Will output assembly as well as the artifact if appropriate for the artifact type for
diff --git a/source/slang/slang-diagnostic-defs.h b/source/slang/slang-diagnostic-defs.h
index f2c7fecc1..21bf73d6e 100644
--- a/source/slang/slang-diagnostic-defs.h
+++ b/source/slang/slang-diagnostic-defs.h
@@ -2702,4 +2702,14 @@ DIAGNOSTIC(
     noBlocksOrIntrinsic,
     "no blocks found for function definition, is there a '$0' intrinsic missing?")
 
+//
+// Ray tracing
+//
+
+DIAGNOSTIC(
+    40000,
+    Error,
+    rayPayloadFieldMissingAccessQualifiers,
+    "field '$0' in ray payload struct must have either 'read' OR 'write' access qualifiers")
+
 #undef DIAGNOSTIC
diff --git a/source/slang/slang-emit-hlsl.cpp b/source/slang/slang-emit-hlsl.cpp
index 0f1ef3ee0..2d963866d 100644
--- a/source/slang/slang-emit-hlsl.cpp
+++ b/source/slang/slang-emit-hlsl.cpp
@@ -1667,8 +1667,18 @@ void HLSLSourceEmitter::emitPostKeywordTypeAttributesImpl(IRInst* inst)
     {
         m_writer->emit("[payload] ");
     }
-    // This can be re-enabled when we add PAQs: https://github.com/shader-slang/slang/issues/3448
-    const bool enablePAQs = false;
+
+    // Get the target profile to determine if PAQs are supported
+    bool enablePAQs = false;
+    auto profile = getTargetProgram()->getOptionSet().getProfile();
+    if (profile.getFamily() == ProfileFamily::DX)
+    {
+        // PAQs are default in Shader Model 6.7 and above when called with `--profile lib_6_7`
+
+        auto version = profile.getVersion();
+        enablePAQs = version >= ProfileVersion::DX_6_7;
+    }
+
     if (enablePAQs)
     {
         if (const auto payloadDecoration = inst->findDecoration<IRRayPayloadDecoration>())
diff --git a/source/slang/slang-lower-to-ir.cpp b/source/slang/slang-lower-to-ir.cpp
index e6ec68660..260596dc3 100644
--- a/source/slang/slang-lower-to-ir.cpp
+++ b/source/slang/slang-lower-to-ir.cpp
@@ -9318,6 +9318,11 @@ struct DeclLoweringVisitor : DeclVisitor<DeclLoweringVisitor, LoweredValInfo>
             subBuilder->addDecoration(irAggType, kIROp_PayloadDecoration);
         }
 
+        if (const auto rayPayloadAttribute = decl->findModifier<RayPayloadAttribute>())
+        {
+            subBuilder->addDecoration(irAggType, kIROp_RayPayloadDecoration);
+        }
+
         subBuilder->setInsertInto(irAggType);
 
         // A `struct` that inherits from another `struct` must start
diff --git a/tests/diagnostics/raypayload-missing-access-qualifiers.slang b/tests/diagnostics/raypayload-missing-access-qualifiers.slang
new file mode 100644
index 000000000..d22a6300b
--- /dev/null
+++ b/tests/diagnostics/raypayload-missing-access-qualifiers.slang
@@ -0,0 +1,37 @@
+// raypayload-missing-access-qualifiers.slang
+
+//DIAGNOSTIC_TEST:SIMPLE:
+
+// Test error for field in ray payload struct missing read/write access qualifiers
+
+struct [raypayload] RayPayload
+{
+    float4 color : read(caller, anyhit) : write(caller);
+    float4 colorMissingQualifiers; // Error expected here
+
+};
+
+uniform RWTexture2D<float4> resultTexture;
+uniform RaytracingAccelerationStructure sceneBVH;
+
+[shader("raygeneration")]
+void rayGenShaderA()
+{
+    int2 threadIdx = DispatchRaysIndex().xy;
+
+    float3 rayDir = float3(0, 0, 1);
+    float3 rayOrigin = 0;
+    rayOrigin.x = (threadIdx.x * 2) - 1;
+    rayOrigin.y = (threadIdx.y * 2) - 1;
+
+    // Trace the ray.
+    RayDesc ray;
+    ray.Origin = rayOrigin;
+    ray.Direction = rayDir;
+    ray.TMin = 0.001;
+    ray.TMax = 10000.0;
+    RayPayload payload = { float4(0, 0, 0, 0) , float4(0, 0, 0, 0)};
+    TraceRay(sceneBVH, RAY_FLAG_NONE, ~0, 0, 0, 0, ray, payload);
+
+    resultTexture[threadIdx.xy] = payload.color;
+}
diff --git a/tests/diagnostics/raypayload-missing-access-qualifiers.slang.expected b/tests/diagnostics/raypayload-missing-access-qualifiers.slang.expected
new file mode 100644
index 000000000..525e8529c
--- /dev/null
+++ b/tests/diagnostics/raypayload-missing-access-qualifiers.slang.expected
@@ -0,0 +1,8 @@
+result code = -1
+standard error = {
+tests/diagnostics/raypayload-missing-access-qualifiers.slang(10): error 40000: field 'colorMissingQualifiers' in ray payload struct must have either 'read' OR 'write' access qualifiers
+    float4 colorMissingQualifiers; // Error expected here
+           ^~~~~~~~~~~~~~~~~~~~~~
+}
+standard output = {
+}
diff --git a/tests/hlsl/raypayload-attribute-no-struct.slang b/tests/hlsl/raypayload-attribute-no-struct.slang
index c7ad94593..4e4921e14 100644
--- a/tests/hlsl/raypayload-attribute-no-struct.slang
+++ b/tests/hlsl/raypayload-attribute-no-struct.slang
@@ -1,7 +1,8 @@
-//enable when https://github.com/shader-slang/slang/issues/3448 is implemented
-//DISABLE_TEST:SIMPLE(filecheck=CHECK): -target hlsl -stage raygeneration -entry rayGenShaderA
+//TEST:SIMPLE(filecheck=CHECK): -target hlsl -profile lib_6_6 -stage raygeneration -entry rayGenShaderA
+//TEST:SIMPLE(filecheck=DXIL): -target dxil -profile lib_6_6 -stage raygeneration -entry rayGenShaderA
 
-// CHECK: struct [raypayload]
+// CHECK: struct RayPayload
+// DXIL: define void @
 
 uniform RWTexture2D resultTexture;
 uniform RaytracingAccelerationStructure sceneBVH;
diff --git a/tests/hlsl/raypayload-attribute-paq.slang b/tests/hlsl/raypayload-attribute-paq.slang
new file mode 100644
index 000000000..3af0556bc
--- /dev/null
+++ b/tests/hlsl/raypayload-attribute-paq.slang
@@ -0,0 +1,37 @@
+//TEST:SIMPLE(filecheck=CHECK): -target hlsl -profile lib_6_7 -stage raygeneration -entry rayGenShaderA
+//TEST:SIMPLE(filecheck=DXIL): -target dxil -profile lib_6_7 -stage raygeneration -entry rayGenShaderA
+
+// CHECK: struct [raypayload]
+// CHECK: float4 color_0 : read(caller, anyhit) : write(caller);
+// DXIL: define void @
+// DXIL: !dx.dxrPayloadAnnotations
+
+struct [raypayload] RayPayload
+{
+    float4 color : read(caller, anyhit) : write(caller);
+};
+
+uniform RWTexture2D<float4> resultTexture;
+uniform RaytracingAccelerationStructure sceneBVH;
+
+[shader("raygeneration")]
+void rayGenShaderA()
+{
+    int2 threadIdx = DispatchRaysIndex().xy;
+
+    float3 rayDir = float3(0, 0, 1);
+    float3 rayOrigin = 0;
+    rayOrigin.x = (threadIdx.x * 2) - 1;
+    rayOrigin.y = (threadIdx.y * 2) - 1;
+
+    // Trace the ray.
+    RayDesc ray;
+    ray.Origin = rayOrigin;
+    ray.Direction = rayDir;
+    ray.TMin = 0.001;
+    ray.TMax = 10000.0;
+    RayPayload payload = { float4(0, 0, 0, 0) };
+    TraceRay(sceneBVH, RAY_FLAG_NONE, ~0, 0, 0, 0, ray, payload);
+
+    resultTexture[threadIdx.xy] = payload.color;
+}
diff --git a/tests/hlsl/raypayload-attribute.slang b/tests/hlsl/raypayload-attribute.slang
index b981589ac..1a9e9a7f5 100644
--- a/tests/hlsl/raypayload-attribute.slang
+++ b/tests/hlsl/raypayload-attribute.slang
@@ -1,8 +1,9 @@
-//enable when https://github.com/shader-slang/slang/issues/3448 is implemented
-//DISABLE_TEST:SIMPLE(filecheck=CHECK): -target hlsl -stage raygeneration -entry rayGenShaderA
-
-// CHECK: struct [raypayload]
+//TEST:SIMPLE(filecheck=CHECK): -target hlsl -profile lib_6_6 -stage raygeneration -entry rayGenShaderA
+//TEST:SIMPLE(filecheck=DXIL): -target dxil -profile lib_6_6 -stage raygeneration -entry rayGenShaderA
 
+// CHECK: struct RayPayload
+// CHECK: float4 color
+// DXIL: define void @
 struct RayPayload
 {
     float4 color;
diff --git a/tools/gfx-unit-test/ray-tracing-test-shaders.slang b/tools/gfx-unit-test/ray-tracing-test-shaders.slang
index aa2e5055f..c1273a717 100644
--- a/tools/gfx-unit-test/ray-tracing-test-shaders.slang
+++ b/tools/gfx-unit-test/ray-tracing-test-shaders.slang
@@ -1,8 +1,8 @@
 // ray-tracing-test-shaders.slang
 
-struct RayPayload
+struct [raypayload] RayPayload
 {
-    float4 color;
+    float4 color : read(caller) : write(caller, closesthit, miss);
 };
 
 uniform RWTexture2D resultTexture;