yum-mirror/slang

Making it easier to work with shaders

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

ArielG-NVInitial copy elision pass (#8042)063cbeaae

master
branchesmasterall refs →
1.7 KiB64 linesraw
1// pointer-self-reference.slang
2
3//TEST:SIMPLE(filecheck=CHECK): -target cuda -entry computeMain -stage compute
4//TEST(compute):COMPARE_COMPUTE_EX(filecheck-buffer=BUFFER): -slang -compute -output-using-type -shaderobj
5//TEST(compute):COMPARE_COMPUTE_EX(filecheck-buffer=BUFFER): -vk -compute -output-using-type -shaderobj
6//TEST(compute):COMPARE_COMPUTE_EX(filecheck-buffer=BUFFER): -cuda -compute -output-using-type -shaderobj
7
8//TEST_INPUT:ubuffer(data=[0 0 0 0], stride=4):out,name=outputBuffer
9RWStructuredBuffer<int> outputBuffer;
10
11struct Thing
12{
13    int value;
14    int bigArray[128];
15
16    // Check that we are not inserting local variables that are copies of `this` parameter.
17
18    // CHECK: __device__ int Thing_getSum{{.*}}Thing{{.*}}*{{.*}}this{{.*}})
19    // CHECK-NOT: Thing{{[a-zA-Z0-9_]*}} {{[a-zA-Z0-9_]+}}
20    // CHECK: }
21    [constref]
22    int getSum()
23    {
24        int result = 0;
25        for (int i = 0; i < 128; i++)
26        {
27            result += bigArray[i];
28        }
29        return result;
30    }
31};
32
33// Check that we are not inserting local variables that are copies of `thing` parameter.
34
35// CHECK: __device__ int test{{.*}}Thing{{.*}}*{{.*}}thing{{.*}})
36// CHECK-NOT: Thing{{[a-zA-Z0-9_]*}} {{[a-zA-Z0-9_]+}}
37// CHECK: }
38
39int test(__constref Thing thing)
40{
41    int sum0 = thing.getSum();
42    AllMemoryBarrier();
43    int sum1 = thing.getSum();
44    return sum0 + sum1;
45}
46
47int caller(Thing thing)
48{
49    return test(thing);
50}
51
52[numthreads(1, 1, 1)]
53void computeMain(int3 dispatchThreadID: SV_DispatchThreadID)
54{
55    int idx = dispatchThreadID.x;
56
57    Thing thing = {};
58
59    thing.bigArray[0] = 100;
60    thing.bigArray[99] = 1;
61
62    // BUFFER: 202
63    outputBuffer[idx] = caller(thing);
64}