Support for more Ray Tracing SPIR-V

btipling · btipling · commit d9566d2a250f · 2025-08-16T19:56:44.000-07:00
diff --git a/include/shady/grammar.json b/include/shady/grammar.json
@@ -73,6 +73,18 @@
       "name": "Image",
       "llvm-id": 397
     },
+    {
+      "name": "IncomingRayPayload",
+      "llvm-id": 399
+    },
+    {
+      "name": "RayPayload",
+      "llvm-id": 400
+    },
+    {
+      "name": "HitAttribute",
+      "llvm-id": 401
+    },
     {
       "name": "CallableDataKHR"
     },
@@ -783,4 +795,4 @@
       ]
     }
   ]
-}
+}
diff --git a/include/shady/ir/execution_model.h b/include/shady/ir/execution_model.h
@@ -9,6 +9,8 @@ EM(Fragment     ) \
 EM(Vertex       ) \
 EM(RayGeneration) \
 EM(Callable     ) \
+EM(RayMiss      ) \
+EM(RayCHit      ) \
 
 typedef enum {
     ShdExecutionModelNone,
@@ -25,6 +27,8 @@ bool shd_get_workgroup_size_for_entry_point(const Node* decl, uint32_t* out);
 static inline bool shd_is_rt_execution_model(ShdExecutionModel em) {
     switch (em) {
         case ShdExecutionModelRayGeneration:
+        case ShdExecutionModelRayMiss:
+        case ShdExecutionModelRayCHit:
         case ShdExecutionModelCallable: return true;
         default: return false;
     }
diff --git a/src/backend/spirv/emit_spv.c b/src/backend/spirv/emit_spv.c
@@ -1,6 +1,7 @@
 #include "emit_spv.h"
 
 #include "shady/ir/builtin.h"
+#include "shady/ir/execution_model.h"
 #include "shady/pass.h"
 
 #include "../shady/ir_private.h"
@@ -257,6 +258,14 @@ void shd_spv_emit_debuginfo(Emitter* emitter, const Node* n, SpvId id) {
 
 static SpvExecutionModel emit_exec_model(Emitter* emitter, ShdExecutionModel model) {
     switch (model) {
+        case ShdExecutionModelRayCHit:
+            spvb_extension(emitter->file_builder, "SPV_KHR_ray_tracing");
+            spvb_capability(emitter->file_builder, SpvCapabilityRayTracingKHR);
+            return SpvExecutionModelClosestHitKHR;
+        case ShdExecutionModelRayMiss:
+            spvb_extension(emitter->file_builder, "SPV_KHR_ray_tracing");
+            spvb_capability(emitter->file_builder, SpvCapabilityRayTracingKHR);
+            return SpvExecutionModelMissKHR;
         case ShdExecutionModelRayGeneration:
             spvb_extension(emitter->file_builder, "SPV_KHR_ray_tracing");
             spvb_capability(emitter->file_builder, SpvCapabilityRayTracingKHR);
diff --git a/src/backend/spirv/emit_spv_type.c b/src/backend/spirv/emit_spv_type.c
@@ -1,12 +1,14 @@
 #include "emit_spv.h"
 
+#include "shady/ir/enum.h"
 #include "shady/ir/memory_layout.h"
 
 #include "shady/rewrite.h"
 
 #include "portability.h"
 #include "log.h"
 #include "dict.h"
+#include "spirv/unified1/spirv.h"
 
 #include <assert.h>
 
@@ -32,6 +34,9 @@ SpvStorageClass spv_emit_addr_space(Emitter* emitter, AddressSpace address_space
         case AsUniformConstant:              return SpvStorageClassUniformConstant;
         case AsIncomingCallableDataKHR:      return SpvStorageClassIncomingCallableDataKHR;
         case AsCallableDataKHR:              return SpvStorageClassCallableDataKHR;
+        case AsIncomingRayPayload:           return SpvStorageClassIncomingRayPayloadKHR;
+        case AsRayPayload:                   return SpvStorageClassRayPayloadKHR;
+        case AsHitAttribute:                 return SpvStorageClassHitAttributeKHR;
 
         case AsCode: return SpvStorageClassCodeSectionSHADY;
 
diff --git a/src/frontend/llvm/l2s_annotations.c b/src/frontend/llvm/l2s_annotations.c
@@ -54,6 +54,9 @@ static bool is_io_as(AddressSpace as) {
         case AsUInput:
         case AsOutput:
         case AsUniform:
+        case AsIncomingRayPayload:
+        case AsRayPayload:
+        case AsHitAttribute:
         case AsUniformConstant: return true;
         default: break;
     }
diff --git a/src/frontend/spirv/s2s.c b/src/frontend/spirv/s2s.c
@@ -892,6 +892,8 @@ static size_t parse_spv_instruction_at(SpvParser* parser, size_t instruction_off
                                            shd_int32_literal(parser->arena, wg_size_dec->payload.literals.data[1]),
                                            shd_int32_literal(parser->arena, wg_size_dec->payload.literals.data[2]))
                     }));
+                } else if (strcmp(entry_point_type->payload.str, "RayCHit") == 0) {
+                } else if (strcmp(entry_point_type->payload.str, "RayMiss") == 0) {
                 } else if (strcmp(entry_point_type->payload.str, "Fragment") == 0) {
 
                 } else if (strcmp(entry_point_type->payload.str, "Vertex") == 0) {
diff --git a/src/shady/passes/io/promote_io_variables.c b/src/shady/passes/io/promote_io_variables.c
@@ -1,4 +1,5 @@
 #include "shady/ir/builtin.h"
+#include "shady/ir/enum.h"
 #include "shady/ir/function.h"
 #include "shady/ir/mem.h"
 #include "shady/ir/debug.h"
@@ -32,7 +33,10 @@ static const Node* promote_to_physical(Context* ctx, AddressSpace as, const Node
         case AsPushConstant:
         case AsUniformConstant:
         case AsUInput:
-        case AsInput: {
+        case AsInput:
+        case AsIncomingRayPayload:
+        case AsRayPayload:
+        case AsHitAttribute: {
             const Node* value = shd_bld_load(ctx->init_bld, io);
             shd_bld_store(ctx->init_bld, phy, value);
             // shd_bld_add_instruction(ctx->init_bld, copy_bytes_helper(a, shd_bld_mem(ctx->init_bld), phy, io, ))
diff --git a/test/vcc/CMakeLists.txt b/test/vcc/CMakeLists.txt
@@ -28,6 +28,9 @@ shady_unit_test(COMPILER vcc NAME vcc_flat_decoration_frag FILES flat_decoration
 shady_unit_test(COMPILER vcc NAME vcc_checker_frag FILES checkerboard.frag.c ARGS ${VCC_TEST_ARGS} --entry-point main --execution-model Fragment VALSPV SPV_VAL_ARGS ${SHADY_SPV_VAL_VK_ARGS})
 shady_unit_test(COMPILER vcc NAME vcc_textured_frag FILES textured.frag.c ARGS ${VCC_TEST_ARGS} --entry-point main --execution-model Fragment VALSPV SPV_VAL_ARGS ${SHADY_SPV_VAL_VK_ARGS})
 
+shady_unit_test(COMPILER vcc NAME vcc_basic_rmiss FILES basic.rmiss.c ARGS ${VCC_TEST_ARGS} --entry-point main --execution-model Fragment VALSPV SPV_VAL_ARGS ${SHADY_SPV_VAL_VK_ARGS})
+shady_unit_test(COMPILER vcc NAME vcc_basic_chit FILES basic.rchit.c ARGS ${VCC_TEST_ARGS} --entry-point main --execution-model Fragment VALSPV SPV_VAL_ARGS ${SHADY_SPV_VAL_VK_ARGS})
+
 shady_unit_test(COMPILER vcc NAME vcc_linkage FILES linkage.frag.c lib.c ARGS ${VCC_TEST_ARGS} --entry-point main --execution-model Fragment VALSPV SPV_VAL_ARGS ${SHADY_SPV_VAL_VK_ARGS})
 shady_unit_test(COMPILER vcc NAME vcc_linkage_type FILES linkage2.frag.c lib2.c ARGS ${VCC_TEST_ARGS} --entry-point main --execution-model Fragment VALSPV SPV_VAL_ARGS ${SHADY_SPV_VAL_VK_ARGS})
 
diff --git a/test/vcc/basic.rchit.c b/test/vcc/basic.rchit.c
@@ -0,0 +1,16 @@
+#include <shady.h>
+
+#define ray_chit_shader  __attribute__((annotate("shady::entry_point::RayCHit")))
+#define ray_payload __attribute__((annotate("shady::io::400")))
+#define hit_attribute __attribute__((annotate("shady::io::401")))
+
+location(0) ray_payload native_vec3 hitValue;
+
+hit_attribute native_vec2 attribs;
+
+ray_chit_shader void main() {
+    const native_vec3 barycentricCoords = (native_vec3){1.f - attribs.x - attribs.y, attribs.x, attribs.y};
+    hitValue = barycentricCoords;
+}
+
+
diff --git a/test/vcc/basic.rmiss.c b/test/vcc/basic.rmiss.c
@@ -0,0 +1,11 @@
+#include <shady.h>
+
+#define ray_miss_shader  __attribute__((annotate("shady::entry_point::RayMiss")))
+#define incoming_ray_payload __attribute__((annotate("shady::io::399")))
+
+location(0) incoming_ray_payload native_vec3 hitValue;
+
+ray_miss_shader void main() {
+    hitValue = (native_vec3){0.f, 0.f, 0.2f};
+}
+

Original file line number	Diff line number	Diff line change
`@@ -73,6 +73,18 @@`
`73`	`73`	`"name": "Image",`
`74`	`74`	`"llvm-id": 397`
`75`	`75`	`},`
	`76`	`+ {`
	`77`	`+ "name": "IncomingRayPayload",`
	`78`	`+ "llvm-id": 399`
	`79`	`+ },`
	`80`	`+ {`
	`81`	`+ "name": "RayPayload",`
	`82`	`+ "llvm-id": 400`
	`83`	`+ },`
	`84`	`+ {`
	`85`	`+ "name": "HitAttribute",`
	`86`	`+ "llvm-id": 401`
	`87`	`+ },`
`76`	`88`	`{`
`77`	`89`	`"name": "CallableDataKHR"`
`78`	`90`	`},`
`@@ -783,4 +795,4 @@`
`783`	`795`	`]`
`784`	`796`	`}`
`785`	`797`	`]`
`786`		`-}`
	`798`	`+}`