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 →
993 B34 linesraw
1//TEST:SIMPLE(filecheck=CUDA): -stage compute -entry computeMain -target cuda -line-directive-mode none
2//TEST:SIMPLE(filecheck=PTX): -stage compute -entry computeMain -target cuda
3
4struct Data {
5    StructuredBuffer<float> input[2];
6    RWStructuredBuffer<float> output;
7    uint input_tensor_count;
8    StructuredBuffer<uint> index_buffer;
9    uint index_count;
10
11    // CUDA: __device__ float Data_fetch{{.*}}(int {{.*}}, int {{.*}})
12    // CUDA-NEXT: {
13    // CUDA-NEXT: globalParams{{.*}}->data{{.*}}->input{{.*}}[{{.*}}]
14    float fetch(int buffer, int index)
15    {
16        return input[buffer][index];
17    }
18};
19
20ParameterBlock<Data> data;
21
22// PTX: computeMain
23
24[shader("compute")]
25[numthreads(8, 8, 1)]
26void computeMain(uint3 tid: SV_DispatchThreadID)
27{
28    float result = 0.0;
29    for (int i = 0; i < data.index_count; ++i) {
30        uint buffer = data.index_buffer[i];
31        result += data.fetch(buffer, tid.x * 1024 + tid.y);
32    }
33    data.output[tid.x * 1024 + tid.y] = result;
34}