yum-mirror/slang

Making it easier to work with shaders

git clone https://git.yummers.dev/yum-mirror/slang

Yong HeImmutable access qualifier for pointers and use `__ldg` on cuda. (#8710)01510f2c9

master
branchesmasterall refs →
1.3 KiB38 linesraw
1//TEST:SIMPLE(filecheck=CUDA): -target cuda -entry compute_main -stage compute
2//TEST:SIMPLE(filecheck=SPV): -target spirv
3
4// Test that we can defer loading big structured buffer elements.
5
6struct Bottom
7{
8    float bigArray[1024];
9    float bottomGetValue(int index) { return bigArray[index]; }
10}
11
12struct Root
13{
14    Bottom bottom;
15}
16
17StructuredBuffer<Root> sb;
18
19RWStructuredBuffer<float> outputBuffer;
20
21// Check that we don't load the entire `Root` struct and then do ElementExtract to get to `bigArray[0]`.
22// Instead we use access chain all the way to point to the required array element, and load just a single float.
23
24// SPV: OpEntryPoint
25// SPV: %[[SBPTRARRAY:[A-Za-z0-9_]+]] = OpAccessChain %_ptr_StorageBuffer__arr_float_int_1024
26// SPV: %[[SBPTR:[A-Za-z0-9_]+]] = OpAccessChain %_ptr_StorageBuffer_float %[[SBPTRARRAY]]
27// SPV: %[[VALUE:[A-Za-z0-9_]+]] = OpLoad %float %[[SBPTR]]
28// SPV: OpStore %{{.*}} %[[VALUE]]
29
30// CUDA: __device__ float Bottom_bottomGetValue{{.*}}(uint [[PARAM0:[A-Za-z0-9_]+]], int [[PARAM1:[A-Za-z0-9_]+]])
31// CUDA:  __ldg(&(&(&(globalParams_0->sb_0){{\[}}[[PARAM0]]{{\]}})->bottom_0)->bigArray_0{{\[}}[[PARAM1]]{{\]}});
32
33[shader("compute")]
34[numthreads(1, 1, 1)]
35void compute_main(uint3 tid: SV_DispatchThreadID)
36{
37    outputBuffer[0] = sb[tid.x].bottom.bottomGetValue(0);
38}