yum-mirror/slang

Making it easier to work with shaders

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

Harsh Aggarwal (NVIDIA)Fix#8086: Batch-10: Enable cuda tests (#8270)639978008

master
branchesmasterall refs →
1.8 KiB55 linesraw
1//TEST:SIMPLE(filecheck=CHECK_SPIRV): -target spirv -entry computeMain -profile cs_6_2 -emit-spirv-via-glsl
2//TEST:SIMPLE(filecheck=CHECK_SPIRV): -target spirv -entry computeMain -profile cs_6_2 -emit-spirv-directly
3//TEST:CROSS_COMPILE: -target dxil-assembly -entry computeMain -profile cs_6_2
4//TEST:SIMPLE(filecheck=CHECK_CUDA): -target cuda -entry computeMain -profile cs_6_2
5
6//TEST_INPUT:ubuffer(data=[0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0], stride=16):out
7RWStructuredBuffer<int> outputBuffer;
8
9//TEST_INPUT: Texture2D(size=4):
10RWTexture2D<half> halfTexture;
11//TEST_INPUT: Texture2D(size=4):
12RWTexture2D<half2> halfTexture2;
13//TEST_INPUT: Texture2D(size=4):
14RWTexture2D<half4> halfTexture4;
15
16//TEST_INPUT: Sampler:
17SamplerState 	s;
18
19[numthreads(4, 4, 1)]
20void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID)
21{
22    int2 pos = int2(dispatchThreadID.xy);
23    float2 uv = pos * (1.0f / 3.0f);
24    int2 pos2 = int2(3 - pos.y, 3 - pos.x);
25    
26#if 0 
27    half  h = halfTexture.Sample(s, uv);
28    half2 h2 = halfTexture2.Sample(s, uv);
29    half4 h4 = halfTexture4.Sample(s, uv);
30#else
31    // CHECK_SPIRV: {{.*}} = OpImageRead
32    // CHECK_SPIRV: {{.*}} = OpFConvert
33    half h = halfTexture[pos2];
34    // CHECK_SPIRV: {{.*}} = OpImageRead
35    // CHECK_SPIRV: {{.*}} = OpFConvert
36    half2 h2 = halfTexture2[pos2];
37    // CHECK_SPIRV: {{.*}} = OpImageRead
38    // CHECK_SPIRV: {{.*}} = OpFConvert
39    half4 h4 = halfTexture4[pos2];
40#endif
41    // Store a results
42    halfTexture[pos] = h2.x + h2.y;
43    halfTexture2[pos] = h4.xy;
44    halfTexture4[pos] = half4(h2, h, h);
45    
46    int index = pos.x + pos.y * 4;
47    outputBuffer[index] = index;
48}
49
50// CHECK_CUDA: surf2Dread<__half>
51// CHECK_CUDA: surf2Dread<__half2{{.*}}>
52// CHECK_CUDA: surf2Dread<__half4{{.*}}>
53// CHECK_CUDA: surf2Dwrite<__half>
54// CHECK_CUDA: surf2Dwrite<__half2{{.*}}>
55// CHECK_CUDA: surf2Dwrite<__half4{{.*}}>