4 files changed, 26 insertions, 3 deletions
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<float> 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<int> [[BUFFER0:b0[_0-9]*]];
+//CHECK_CUDA-DAG: RWStructuredBuffer<int> [[BUFFER1:b1[_0-9]*]];
+//CHECK_CUDA-DAG: RWStructuredBuffer<int> [[BUFFER2:b2[_0-9]*]];
+//CHECK_CUDA-DAG: RWStructuredBuffer<int> [[BUFFER3:b3[_0-9]*]];
+//CHECK_CUDA-DAG: FixedArray<RWStructuredBuffer<int>, 2>  [[BUFFER4:b4[_0-9]*]];
+//CHECK_CUDA-DAG: FixedArray<RWStructuredBuffer<int>, 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<int> 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<int> 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<int> [[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)