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/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 +++++++++-- 4 files changed, 26 insertions(+), 3 deletions(-) (limited to 'tests/hlsl') 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