From 639978008de3a74c00e03451cd9fc74452766fcd Mon Sep 17 00:00:00 2001 From: "Harsh Aggarwal (NVIDIA)" Date: Wed, 3 Sep 2025 10:45:57 +0530 Subject: Fix#8086: Batch-10: Enable cuda tests (#8270) --- tests/compute/half-texture.slang | 8 ++++++++ tests/compute/mutating-and-inout.slang | 3 ++- tests/compute/typedef-member.slang | 3 ++- tests/hlsl/byte-buffer-load-std430.slang | 1 + tests/hlsl/register-syntax.slang | 15 +++++++++++++++ tests/hlsl/texture-mips-operator.slang | 2 +- tests/hlsl/wave-size.slang | 11 +++++++++-- 7 files changed, 38 insertions(+), 5 deletions(-) diff --git a/tests/compute/half-texture.slang b/tests/compute/half-texture.slang index ca51df3cc..6cda1f4c9 100644 --- a/tests/compute/half-texture.slang +++ b/tests/compute/half-texture.slang @@ -1,6 +1,7 @@ //TEST:SIMPLE(filecheck=CHECK_SPIRV): -target spirv -entry computeMain -profile cs_6_2 -emit-spirv-via-glsl //TEST:SIMPLE(filecheck=CHECK_SPIRV): -target spirv -entry computeMain -profile cs_6_2 -emit-spirv-directly //TEST:CROSS_COMPILE: -target dxil-assembly -entry computeMain -profile cs_6_2 +//TEST:SIMPLE(filecheck=CHECK_CUDA): -target cuda -entry computeMain -profile cs_6_2 //TEST_INPUT:ubuffer(data=[0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0], stride=16):out RWStructuredBuffer outputBuffer; @@ -45,3 +46,10 @@ void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID) int index = pos.x + pos.y * 4; outputBuffer[index] = index; } + +// CHECK_CUDA: surf2Dread<__half> +// CHECK_CUDA: surf2Dread<__half2{{.*}}> +// CHECK_CUDA: surf2Dread<__half4{{.*}}> +// CHECK_CUDA: surf2Dwrite<__half> +// CHECK_CUDA: surf2Dwrite<__half2{{.*}}> +// CHECK_CUDA: surf2Dwrite<__half4{{.*}}> diff --git a/tests/compute/mutating-and-inout.slang b/tests/compute/mutating-and-inout.slang index 37cc30870..11132e407 100644 --- a/tests/compute/mutating-and-inout.slang +++ b/tests/compute/mutating-and-inout.slang @@ -2,7 +2,8 @@ // Test that calling a `[mutating]` method on an `inout` function parameter works. -//TEST(compute):COMPARE_COMPUTE: -shaderobj +//TEST(compute):COMPARE_COMPUTE: -cuda +//TEST(compute):COMPARE_COMPUTE: -vk //TEST_INPUT:ubuffer(data=[0 0 0 0], stride=4):out,name=outputBuffer RWStructuredBuffer outputBuffer; diff --git a/tests/compute/typedef-member.slang b/tests/compute/typedef-member.slang index ae3debd40..98694c839 100644 --- a/tests/compute/typedef-member.slang +++ b/tests/compute/typedef-member.slang @@ -1,4 +1,5 @@ //TEST(compute):COMPARE_COMPUTE: -shaderobj +//TEST(compute):COMPARE_COMPUTE: -shaderobj -cuda // Confirm that a struct type defined in a generic parent works @@ -35,4 +36,4 @@ void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID) float inVal = float(tid); float outVal = test(inVal); outputBuffer[tid] = outVal.x; -} \ No newline at end of file +} diff --git a/tests/hlsl/byte-buffer-load-std430.slang b/tests/hlsl/byte-buffer-load-std430.slang index b526250eb..1e495d310 100644 --- a/tests/hlsl/byte-buffer-load-std430.slang +++ b/tests/hlsl/byte-buffer-load-std430.slang @@ -1,4 +1,5 @@ //TEST(compute, vulkan):COMPARE_COMPUTE_EX:-vk -compute -shaderobj -output-using-type -xslang -fvk-use-gl-layout +//TEST(compute, vulkan):COMPARE_COMPUTE_EX:-cuda -compute -shaderobj -output-using-type -xslang -fvk-use-gl-layout //TEST_INPUT:ubuffer(data=[0], stride=4):out,name=outputBuffer RWStructuredBuffer outputBuffer; diff --git a/tests/hlsl/register-syntax.slang b/tests/hlsl/register-syntax.slang index f3633028d..6eb1241b9 100644 --- a/tests/hlsl/register-syntax.slang +++ b/tests/hlsl/register-syntax.slang @@ -1,5 +1,6 @@ //TEST:SIMPLE(filecheck=CHECK_GLSL): -target glsl -stage compute -entry computeMain //TEST:SIMPLE(filecheck=CHECK_HLSL): -target hlsl -stage compute -entry computeMain +//TEST:SIMPLE(filecheck=CHECK_CUDA): -target cuda -stage compute -entry computeMain //CHECK_GLSL-DAG: binding = 4, set = 2 //CHECK_GLSL-DAG: binding = 5, set = 1 @@ -15,6 +16,20 @@ //CHECK_HLSL-DAG: u9 //CHECK_HLSL-DAG: u12 +//CHECK_CUDA-DAG: RWStructuredBuffer [[BUFFER0:b0[_0-9]*]]; +//CHECK_CUDA-DAG: RWStructuredBuffer [[BUFFER1:b1[_0-9]*]]; +//CHECK_CUDA-DAG: RWStructuredBuffer [[BUFFER2:b2[_0-9]*]]; +//CHECK_CUDA-DAG: RWStructuredBuffer [[BUFFER3:b3[_0-9]*]]; +//CHECK_CUDA-DAG: FixedArray, 2> [[BUFFER4:b4[_0-9]*]]; +//CHECK_CUDA-DAG: FixedArray, 2> [[BUFFER5:b5[_0-9]*]]; +//CHECK_CUDA: *(&([[GLOBALPARAMS:globalParams[_0-9]*]]->[[BUFFER0]])[int(0)]) = int(1); +//CHECK_CUDA: *(&([[GLOBALPARAMS]]->[[BUFFER1]])[int(0)]) = int(1); +//CHECK_CUDA: *(&([[GLOBALPARAMS]]->[[BUFFER2]])[int(0)]) = int(1); +//CHECK_CUDA: *(&([[GLOBALPARAMS]]->[[BUFFER3]])[int(0)]) = int(1); +//CHECK_CUDA: *(&([[GLOBALPARAMS]]->[[BUFFER4]][int(0)])[int(0)]) = int(1); +//CHECK_CUDA: *(&([[GLOBALPARAMS]]->[[BUFFER4]][int(0)])[int(1)]) = int(1); +//CHECK_CUDA: *(&([[GLOBALPARAMS]]->[[BUFFER5]][int(0)])[int(1)]) = int(1); + [[vk::binding(4,2)]] RWStructuredBuffer b0 : register(u4, space2); diff --git a/tests/hlsl/texture-mips-operator.slang b/tests/hlsl/texture-mips-operator.slang index 9da8711c5..4b3b7fd98 100644 --- a/tests/hlsl/texture-mips-operator.slang +++ b/tests/hlsl/texture-mips-operator.slang @@ -10,4 +10,4 @@ uniform Texture2D gTex; void main() { result[0] = gTex.mips[1][int2(3,4)]; -} \ No newline at end of file +} diff --git a/tests/hlsl/wave-size.slang b/tests/hlsl/wave-size.slang index 3fc6363c4..4a30aafaa 100644 --- a/tests/hlsl/wave-size.slang +++ b/tests/hlsl/wave-size.slang @@ -1,9 +1,16 @@ -//TEST:SIMPLE(filecheck=CHECK): -target hlsl -stage compute -entry computeMain +//TEST:SIMPLE(filecheck=CHECK_HLSL): -target hlsl -stage compute -entry computeMain +//TEST:SIMPLE(filecheck=CHECK_CUDA): -target cuda -stage compute -entry computeMain //TEST_INPUT:ubuffer(data=[0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0], stride=4):out,name=outputBuffer RWStructuredBuffer outputBuffer; -// CHECK: [WaveSize(4)] +// HLSL preserves WaveSize attribute in generated code +// CHECK_HLSL: [WaveSize(4)] + +// CUDA doesn't emit WaveSize attribute but generates valid compute kernel +// CHECK_CUDA: RWStructuredBuffer [[OUTPUTBUFFER:outputBuffer[_0-9]*]]; +// CHECK_CUDA: extern "C" __global__ void computeMain() +// CHECK_CUDA: *(&([[GLOBALPARAMS:globalParams[_0-9]*]]->[[OUTPUTBUFFER]])[[[TID:tid[_0-9]*]]]) = [[TID]]; [WaveSize(4)] [numthreads(4, 1, 1)] void computeMain(int3 dispatchThreadID : SV_DispatchThreadID) -- cgit v1.2.3