yum-mirror/slang

Making it easier to work with shaders

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

Yong HeEnhance buffer load specialization pass to specialize past field extracts. (#8547)e4611e2e3

master
branchesmasterall refs →
1.4 KiB51 linesraw
1//TEST:SIMPLE(filecheck=CUDA): -target cuda -stage compute -entry compute_main
2//TEST:SIMPLE(filecheck=PTX): -target ptx -stage compute -entry compute_main
3//TEST:SIMPLE(filecheck=SPV): -target spirv
4
5struct Bottom
6{
7    float bigArray[1024];
8
9    [mutating]
10    void setVal(int index, float value) { bigArray[index] = value; }
11}
12
13struct Root
14{
15    Bottom top[2];
16    [mutating]
17    void setTopVal(int x, int y, float value)
18    {
19        top[x].setVal(y, value);
20    }
21}
22
23RWStructuredBuffer<Root> sb;
24
25// Check that we don't load the entire `Root` struct, modify it, and then write it back.
26// Instead we should generate a single store instruction to write the single float value
27// directly to the buffer.
28
29// SPV: OpEntryPoint
30// SPV: OpLabel
31// SPV-NEXT: OpAccessChain
32// SPV-NOT: OpCompositeInsert
33// SPV-NOT: OpLoad
34// SPV: OpStore
35// SPV-NOT: OpLoad
36// SPV-NOT: OpCompositeInsert
37// SPV: OpStore
38// SPV: OpReturn
39
40// CUDA: __device__ void Bottom_setVal_0(int [[INDEX0:[A-Za-z0-9_]+]], int [[INDEX1:[A-Za-z0-9_]+]], int [[INDEX2:[A-Za-z0-9_]+]], float [[VAL:[A-Za-z0-9_]+]])
41// CUDA:  (&(&(globalParams{{.*}}->sb{{.*}}){{\[}}[[INDEX0]]{{\]}})->top{{.*}}{{\[}}[[INDEX1]]{{\]}})->bigArray{{.*}}{{\[}}[[INDEX2]]{{\]}} = [[VAL]];
42// PTX: compute_main
43
44[shader("compute")]
45[numthreads(1, 1, 1)]
46void compute_main(uint3 tid: SV_DispatchThreadID)
47{
48    sb[0].setTopVal(1, 2, 100.0f);
49
50    sb[3].top[1].setVal(8, 200.0f);
51}