Support for Payload Access Qualifiers (#3448)

source/compiler-core/slang-dxc-compiler.cpp

@@ -479,8 +479,47 @@ 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;
+    // According to DirectX Raytracing Specification, PAQs are supported in Shader Model 6.7 and
+    // above
+    bool enablePAQs = false;
+    String profileName = asString(options.profileName);
+    // Check if we're targeting SM 6.7 or higher
+    if (profileName.getLength() > 0)
+    {
+        const UnownedStringSlice profileSlice = profileName.getUnownedSlice();
+        if (profileSlice.startsWith("lib_") || profileSlice.startsWith("sm_"))
+        {
+            // Format is like "lib_6_7" - we want to extract the "6" and "7" parts
+
+            // Get everything after the prefix (e.g., "6_7")
+            const char* start = profileSlice.begin() + (profileSlice.startsWith("lib_") ? 4 : 3);
+            UnownedStringSlice versionString(start, profileSlice.end());
+
+            // Split by '_' to get major and minor
+            UnownedStringSlice majorSlice = StringUtil::getAtInSplit(versionString, '_', 0);
+            UnownedStringSlice minorSlice = StringUtil::getAtInSplit(versionString, '_', 1);
+
+            Int majorVersion = 0;
+            Int minorVersion = 0;
+
+            if (SLANG_SUCCEEDED(StringUtil::parseInt(majorSlice, majorVersion)))
+            {
+                if (majorVersion >= 7)
+                {
+                    // Any shader model 7.x or higher supports PAQs
+                    enablePAQs = true;
+                }
+                else if (
+                    majorVersion == 6 &&
+                    SLANG_SUCCEEDED(StringUtil::parseInt(minorSlice, minorVersion)))
+                {
+                    // For 6.x, we need to check the minor version (must be >= 7)
+                    enablePAQs = minorVersion >= 7;
+                }
+            }
+        }
+    }
+
     if (!enablePAQs)
         args.add(L"-disable-payload-qualifiers");
     else
@@ -560,7 +599,6 @@ SlangResult DXCDownstreamCompiler::compile(const CompileOptions& inOptions, IArt
     //
     args.add(L"-no-warnings");
 
-    String profileName = asString(options.profileName);
     // If we are going to link we have to compile in the lib profile style
     if (libraries.getCount() && hasSource)
     {

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

@@ -12494,6 +12494,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-diagnostic-defs.h

@@ -2694,4 +2694,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

@@ -9207,6 +9207,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/hlsl/raypayload-attribute-no-struct.slang

@@ -1,5 +1,4 @@
-//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_7 -stage raygeneration -entry rayGenShaderA
 
 // CHECK: struct [raypayload]
 

tests/hlsl/raypayload-attribute-paq.slang

@@ -0,0 +1,34 @@
+//TEST:SIMPLE(filecheck=CHECK): -target hlsl -profile lib_6_7 -stage raygeneration -entry rayGenShaderA
+
+// CHECK: struct [raypayload]
+
+struct [raypayload] RayPayload
+{
+    float4 color : write(caller) : read(anyhit);
+
+};
+
+uniform RWTexture2D 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,5 +1,4 @@
-//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_7 -stage raygeneration -entry rayGenShaderA
 
 // CHECK: struct [raypayload]
 

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;