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.6 KiB47 linesraw
1//TEST:SIMPLE(filecheck=CHECK): -target cuda -line-directive-mode none
2
3// This test verifies that DispatchThreadID parameter of different types 
4// correctly extracts components from the underlying uint3 value in CUDA.
5
6//TEST_INPUT: ubuffer(data=[0 0 0 0 0 0 0], stride=4):out,name cudaOutputBuffer
7RWStructuredBuffer<float> cudaOutputBuffer;
8
9[shader("compute")]
10[numthreads(1, 1, 1)]
11void computeMain(uint tid: SV_DispatchThreadID, StructuredBuffer<uint> src, RWStructuredBuffer<uint> dst)
12{
13    dst[tid.x] = src[tid.x];
14}
15// CHECK: uint {{.*}} = (blockIdx * blockDim + threadIdx).x;
16
17[shader("compute")]
18[numthreads(1, 1, 1)]
19void computeMain2(uint2 tid: SV_DispatchThreadID, StructuredBuffer<uint> src, RWStructuredBuffer<uint> dst)
20{
21    dst[tid.x] = src[tid.y];
22}
23// CHECK: uint2 {{.*}} = uint2 {(blockIdx * blockDim + threadIdx).x, (blockIdx * blockDim + threadIdx).y};
24
25[shader("compute")]
26[numthreads(1, 1, 1)]
27void computeMain3(int2 tid: SV_DispatchThreadID, StructuredBuffer<uint> src, RWStructuredBuffer<uint> dst)
28{
29    dst[tid.x] = src[tid.x];
30}
31// CHECK: uint2 {{.*}} = uint2 {(blockIdx * blockDim + threadIdx).x, (blockIdx * blockDim + threadIdx).y};
32
33[shader("compute")]
34[numthreads(1, 1, 1)]
35void computeMain4(int tid: SV_DispatchThreadID, StructuredBuffer<uint> src, RWStructuredBuffer<uint> dst)
36{
37    dst[tid.x] = src[tid.x];
38}
39// CHECK: int {{.*}} = int((blockIdx * blockDim + threadIdx).x);
40
41[shader("compute")]
42[numthreads(1, 1, 1)]
43void computeMain5(int tid: SV_GroupIndex, StructuredBuffer<uint> src, RWStructuredBuffer<uint> dst)
44{
45    dst[tid.x] = src[tid.x];
46}
47// CHECK: int {{.*}} = int((threadIdx.z * blockDim.y + threadIdx.y) * blockDim.x + threadIdx.x);