diff options
Diffstat (limited to 'tests')
| -rw-r--r-- | tests/cuda/dispatch-thread-id-extraction.slang | 9 | ||||
| -rw-r--r-- | tests/language-feature/enums/enum-switch-2.slang | 1 | ||||
| -rw-r--r-- | tests/language-feature/system-value-extraction.slang | 31 |
3 files changed, 37 insertions, 4 deletions
diff --git a/tests/cuda/dispatch-thread-id-extraction.slang b/tests/cuda/dispatch-thread-id-extraction.slang index 5fc3c89a6..02705ff24 100644 --- a/tests/cuda/dispatch-thread-id-extraction.slang +++ b/tests/cuda/dispatch-thread-id-extraction.slang @@ -28,8 +28,9 @@ void computeMain3(int2 tid: SV_DispatchThreadID, StructuredBuffer<uint> src, RWS { dst[tid.x] = src[tid.x]; } -// CHECK: int _S3 = (slang_bit_cast<int2 >(uint2 {(blockIdx * blockDim + threadIdx).x, (blockIdx * blockDim + threadIdx).y})).x; - +// CHECK: uint2 _S3 = uint2 {(blockIdx * blockDim + threadIdx).x, (blockIdx * blockDim + threadIdx).y}; +// CHECK: int2 _S4 = make_int2 ((int)_S3.x, (int)_S3.y); +// CHECK: int _S5 = _S4.x; [shader("compute")] [numthreads(1, 1, 1)] @@ -37,7 +38,7 @@ void computeMain4(int tid: SV_DispatchThreadID, StructuredBuffer<uint> src, RWSt { dst[tid.x] = src[tid.x]; } -// CHECK: int _S4 = (slang_bit_cast<int>((blockIdx * blockDim + threadIdx).x)); +// CHECK: int _S6 = int((blockIdx * blockDim + threadIdx).x); [shader("compute")] [numthreads(1, 1, 1)] @@ -45,4 +46,4 @@ void computeMain5(int tid: SV_GroupIndex, StructuredBuffer<uint> src, RWStructur { dst[tid.x] = src[tid.x]; } -// CHECK: int _S5 = (slang_bit_cast<int>((threadIdx.z * blockDim.y + threadIdx.y) * blockDim.x + threadIdx.x));
\ No newline at end of file +// CHECK: int _S7 = int((threadIdx.z * blockDim.y + threadIdx.y) * blockDim.x + threadIdx.x); diff --git a/tests/language-feature/enums/enum-switch-2.slang b/tests/language-feature/enums/enum-switch-2.slang index f5266f35d..de4a4757b 100644 --- a/tests/language-feature/enums/enum-switch-2.slang +++ b/tests/language-feature/enums/enum-switch-2.slang @@ -1,4 +1,5 @@ //TEST(compute):COMPARE_COMPUTE_EX(filecheck-buffer=CHECK):-slang -compute -shaderobj -output-using-type +//TEST(compute):COMPARE_COMPUTE_EX(filecheck-buffer=CHECK):-cpu -compute -shaderobj -output-using-type //TEST(compute, vulkan):COMPARE_COMPUTE_EX(filecheck-buffer=CHECK):-vk -compute -shaderobj -output-using-type enum class E : uint32_t diff --git a/tests/language-feature/system-value-extraction.slang b/tests/language-feature/system-value-extraction.slang new file mode 100644 index 000000000..12ca07a96 --- /dev/null +++ b/tests/language-feature/system-value-extraction.slang @@ -0,0 +1,31 @@ +//TEST(compute):COMPARE_COMPUTE_EX(filecheck-buffer=CHECK): -dx11 -compute +//TEST(compute):COMPARE_COMPUTE_EX(filecheck-buffer=CHECK): -dx12 -compute +//TEST(compute):COMPARE_COMPUTE_EX(filecheck-buffer=CHECK): -cpu -compute +//TEST(compute):COMPARE_COMPUTE_EX(filecheck-buffer=CHECK): -cuda -compute +//TEST(compute):COMPARE_COMPUTE_EX(filecheck-buffer=CHECK): -wgsl -compute +//DISABLE_TEST(compute):COMPARE_COMPUTE_EX(filecheck-buffer=CHECK): -mtl -compute +//TEST(compute,vulkan):COMPARE_COMPUTE_EX(filecheck-buffer=CHECK): -vk -compute + +// Slang allows the type of system value semantics to differ from their +// "canonical" type - e.g. `uint8_t tid : SV_DispatchThreadID` is technically +// valid and refers to the first element in the underlying uint3 vector. + +// TEST_INPUT: ubuffer(data=[0 0 0 0 0 0 0 0], stride=4):out,name=outputBuffer +RWStructuredBuffer<int> outputBuffer; + +[numthreads(2, 4, 1)] +void computeMain( + uint ind : SV_GroupIndex, + vector<uint, 2> tid: SV_DispatchThreadID, + int localTid : SV_GroupThreadID +){ + // CHECK: 0 + // CHECK-NEXT: 5 + // CHECK-NEXT: 8 + // CHECK-NEXT: D + // CHECK-NEXT: 10 + // CHECK-NEXT: 15 + // CHECK-NEXT: 18 + // CHECK-NEXT: 1D + outputBuffer[ind] = 4 * (tid.x + 2 * tid.y) + localTid; +} |