Support for Payload Access Qualifiers (#3448)

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;

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

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;
+

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

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

source/slang/slang-check-impl.h

@@ -2835,6 +2835,8 @@ struct SemanticsVisitor : public SemanticsContext
     bool isCStyleType(Type* type, HashSet<Type*>& isVisit);
 
     void addVisibilityModifier(Decl* decl, DeclVisibility vis);
+
+    void checkRayPayloadStructFields(StructDecl* structDecl);
 };
 
 

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

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();

source/slang/slang-compiler.h

@@ -2812,7 +2812,9 @@ struct CodeGenContext
     };
 
     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 @@ struct CodeGenContext
 
 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

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

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>())

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

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;
+}

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 = {
+}

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;

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;
+}

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;

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;