Enable CUDA Test Enablement - Batch 1: Autodiff Tests (1-16) (#8139)

11 files changed, 36 insertions, 27 deletions

diff --git a/tests/autodiff/dynamic-dispatch-ptr.slang b/tests/autodiff/dynamic-dispatch-ptr.slang
index 3f2269f78..5a0614769 100644
--- a/tests/autodiff/dynamic-dispatch-ptr.slang
+++ b/tests/autodiff/dynamic-dispatch-ptr.slang
@@ -1,4 +1,5 @@
 //TEST:COMPARE_COMPUTE(filecheck-buffer=CHECK):-vk -output-using-type -emit-spirv-directly
+//TEST:COMPARE_COMPUTE(filecheck-buffer=CHECK):-cuda -output-using-type
 
 //CHECK: 1.0 
 
@@ -40,4 +41,4 @@ void computeMain(
     DifferentialPair<float4> dp;
     bwd_diff(splat)(s, dp, float4(1.0f));
     outBuffer[id.x] = dp.d;
-}
\ No newline at end of file
+}
diff --git a/tests/autodiff/force-inline-differentiable.slang b/tests/autodiff/force-inline-differentiable.slang
index 5d22a525f..a9a2f97a9 100644
--- a/tests/autodiff/force-inline-differentiable.slang
+++ b/tests/autodiff/force-inline-differentiable.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): -slang -compute -shaderobj -output-using-type -cuda
 
 //TEST_INPUT:ubuffer(data=[0 0 0 0 0], stride=4):out,name=outputBuffer
 RWStructuredBuffer<float> outputBuffer;
@@ -44,4 +45,4 @@ void computeMain(uint3 id : SV_DispatchThreadID)
     
     // CHECK: type: float
     // CHECK: 3.0
-}
\ No newline at end of file
+}
diff --git a/tests/autodiff/generic-accessors.slang b/tests/autodiff/generic-accessors.slang
index 2b179f256..ad03c6d79 100644
--- a/tests/autodiff/generic-accessors.slang
+++ b/tests/autodiff/generic-accessors.slang
@@ -1,4 +1,5 @@
 //TEST:COMPARE_COMPUTE(filecheck-buffer=CHK): -output-using-type
+//TEST:COMPARE_COMPUTE(filecheck-buffer=CHK): -output-using-type -cuda
 
 interface ITest
 {
@@ -29,4 +30,4 @@ void computeMain()
     Test t = {};
     output[0] = test(t);
     // CHK: 5.0
-}
\ No newline at end of file
+}
diff --git a/tests/autodiff/property.slang b/tests/autodiff/property.slang
index e15b9a75a..2a626546f 100644
--- a/tests/autodiff/property.slang
+++ b/tests/autodiff/property.slang
@@ -1,4 +1,5 @@
 //TEST:COMPARE_COMPUTE(filecheck-buffer=CHECK):-output-using-type
+//TEST:COMPARE_COMPUTE(filecheck-buffer=CHECK):-output-using-type -cuda
 public struct ReadOnlyIndex
 {
     private int _idx;
@@ -45,4 +46,4 @@ void computeMain()
 {
     // CHECK: 5.0
     output[0] = repro(gPrimal, gGrad);
-}
\ No newline at end of file
+}
diff --git a/tests/autodiff/struct-without-diff-associations.slang b/tests/autodiff/struct-without-diff-associations.slang
index dea7b9630..81f0c9836 100644
--- a/tests/autodiff/struct-without-diff-associations.slang
+++ b/tests/autodiff/struct-without-diff-associations.slang
@@ -1,5 +1,5 @@
 //TEST:SIMPLE(filecheck=CHECK): -target hlsl -line-directive-mode none -entry computeMain -stage compute
-
+//TEST:SIMPLE(filecheck=CHECK): -target cuda -line-directive-mode none -entry computeMain -stage compute
 //TEST_INPUT:ubuffer(data=[0 0 0 0], stride=4):out,name=outputBuffer
 RWStructuredBuffer<float> outputBuffer;
 
@@ -13,7 +13,7 @@ struct Foo : IDifferentiable
     typealias Differential = Foo2;
 
     float x[3];
-    // CHECK: tests/autodiff/struct-without-diff-associations.slang(15): error 30102: differentiable member 'x' should have a corresponding field in 'Foo2'. Use [DerivativeMember(Foo2.<field-name>)] or mark as no_diff
+    // CHECK: tests/autodiff/struct-without-diff-associations.slang([[# @LINE-1]]): error 30102: differentiable member 'x' should have a corresponding field in 'Foo2'. Use [DerivativeMember(Foo2.<field-name>)] or mark as no_diff
     // CHECK-NEXT: float x[3];
     // CHECK-NEXT:       ^
 };
@@ -40,4 +40,4 @@ void computeMain(uint3 dispatchThreadID: SV_DispatchThreadID)
         var d = fwd_diff(foobar)(diffPair(a, 1.0)).d;
         outputBuffer[0] = d;
     }
-}
\ No newline at end of file
+}
diff --git a/tests/autodiff/subscript.slang b/tests/autodiff/subscript.slang
index 2b16597c0..01c171ba0 100644
--- a/tests/autodiff/subscript.slang
+++ b/tests/autodiff/subscript.slang
@@ -1,4 +1,5 @@
 //TEST:COMPARE_COMPUTE(filecheck-buffer=CHK): -output-using-type
+//TEST:COMPARE_COMPUTE(filecheck-buffer=CHK): -output-using-type -cuda
 
 interface ITest
 {
@@ -30,4 +31,4 @@ void computeMain()
     Test t = {};
     output[0] = test(t);
     // CHK: 5.0
-}
\ No newline at end of file
+}
diff --git a/tests/autodiff/test-minimal-context.slang b/tests/autodiff/test-minimal-context.slang
index c2a2b87ed..52296294d 100644
--- a/tests/autodiff/test-minimal-context.slang
+++ b/tests/autodiff/test-minimal-context.slang
@@ -1,4 +1,5 @@
 //TEST:SIMPLE(filecheck=CHECK): -target hlsl -profile cs_5_0 -entry computeMain -line-directive-mode none
+//TEST:SIMPLE(filecheck=CHECK): -target cuda -profile cs_5_0 -entry computeMain -line-directive-mode none
 //DISABLE_TEST:SIMPLE(filecheck=CTX):-target glsl -stage compute -entry computeMain -report-checkpoint-intermediates
 
 //TEST_INPUT:ubuffer(data=[0 0 0 0], stride=4):out,name=outputBuffer
@@ -73,4 +74,4 @@ void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID)
     bwd_diff(f)(0, dpa, 1.0f);
     outputBuffer[0] = dpa.d; // Expect: 1
 }
-// CTX: note:
\ No newline at end of file
+// CTX: note:
diff --git a/tests/autodiff/warn-on-prefer-recompute-side-effects.slang b/tests/autodiff/warn-on-prefer-recompute-side-effects.slang
index b7dc8cf2a..64aaf7666 100644
--- a/tests/autodiff/warn-on-prefer-recompute-side-effects.slang
+++ b/tests/autodiff/warn-on-prefer-recompute-side-effects.slang
@@ -1,4 +1,5 @@
 //TEST:SIMPLE(filecheck=CHECK): -target hlsl -line-directive-mode none
+//TEST:SIMPLE(filecheck=CHECK): -target cuda -line-directive-mode none
 
 //TEST_INPUT:ubuffer(data=[0 0 0 0], stride=4):out,name=outputBuffer
 RWStructuredBuffer<float> outputBuffer;
@@ -12,7 +13,7 @@ float get_thread_5_value(float v, uint group_thread_id)
     if(group_thread_id == 5)
     {
         s_shared = detach(v);
-        // CHECK: tests/autodiff/warn-on-prefer-recompute-side-effects.slang(10): warning 42050: get_thread_5_value has [PreferRecompute] and may have side effects. side effects may execute multiple times. use [PreferRecompute(SideEffectBehavior.Allow)], or mark function with [__NoSideEffect]
+        // CHECK: tests/autodiff/warn-on-prefer-recompute-side-effects.slang([[# @LINE-5]]): warning 42050: get_thread_5_value has [PreferRecompute] and may have side effects. side effects may execute multiple times. use [PreferRecompute(SideEffectBehavior.Allow)], or mark function with [__NoSideEffect]
         // CHECK: float get_thread_5_value(float v, uint group_thread_id)
         // CHECK:       ^~~~~~~~~~~~~~~~~~
     }
@@ -44,4 +45,4 @@ void computeMain(uint3 group_thread_id: SV_GroupThreadID, uint3 dispatch_thread_
     bwd_diff(get_thread_6_value)(value, group_thread_id.x, 1.0f);
 
     outputBuffer[dispatch_thread_id.x] = value.d;
-}
\ No newline at end of file
+}
diff --git a/tests/autodiff/warn-on-shared-memory-access.slang b/tests/autodiff/warn-on-shared-memory-access.slang
index bccf8b1fa..829191d0e 100644
--- a/tests/autodiff/warn-on-shared-memory-access.slang
+++ b/tests/autodiff/warn-on-shared-memory-access.slang
@@ -1,4 +1,5 @@
 //TEST:SIMPLE(filecheck=CHECK): -target hlsl -line-directive-mode none
+//TEST:SIMPLE(filecheck=CHECK): -target cuda -line-directive-mode none
 
 //TEST_INPUT:ubuffer(data=[0 0 0 0], stride=4):out,name=outputBuffer
 RWStructuredBuffer<float> outputBuffer;
@@ -12,7 +13,7 @@ float get_thread_5_value(float v, uint group_thread_id)
     {
         // Using 'detach(v)' makes the error go away
         s_shared = v;
-        // CHECK: tests/autodiff/warn-on-shared-memory-access.slang(14): error 41024: derivative is lost during assignment to non-differentiable location, use 'detach()' to clarify intention.
+        // CHECK: tests/autodiff/warn-on-shared-memory-access.slang([[# @LINE-1]]): error 41024: derivative is lost during assignment to non-differentiable location, use 'detach()' to clarify intention.
         // CHECK: s_shared = v;
         // CHECK:          ^
     }
@@ -29,4 +30,4 @@ void computeMain(uint3 group_thread_id: SV_GroupThreadID, uint3 dispatch_thread_
     bwd_diff(get_thread_5_value)(value, group_thread_id.x, 1.0f);
 
     outputBuffer[dispatch_thread_id.x] = value.d;
-}
\ No newline at end of file
+}
diff --git a/tests/autodiff/was/warped-sampling-1d.slang b/tests/autodiff/was/warped-sampling-1d.slang
index 3a2ca8f92..38e1410fa 100644
--- a/tests/autodiff/was/warped-sampling-1d.slang
+++ b/tests/autodiff/was/warped-sampling-1d.slang
@@ -1,4 +1,5 @@
-//TEST(compute):COMPARE_COMPUTE_EX:-slang -compute -shaderobj -output-using-type -profile cs_5_1 -dx12
+//TEST(compute):COMPARE_COMPUTE_EX(filecheck-buffer=CHECK):-slang -compute -shaderobj -output-using-type -profile cs_5_1 -dx12 -compute-dispatch 4,1,1
+//TEST(compute):COMPARE_COMPUTE_EX(filecheck-buffer=CHECK):-slang -compute -shaderobj -output-using-type -profile cs_5_1 -cuda -compute-dispatch 4,1,1
 
 //TEST_INPUT:ubuffer(data=[0.0 0.0 0.0 0.0 0.0 0.0 0.0 0.0 0.0 0.0], stride=4):out,name=endpointDifferentialBuffer
 RWStructuredBuffer<float> endpointDifferentialBuffer;
@@ -256,7 +257,7 @@ float renderSample(inout PRNG prng)
     return shadeIntersection(isect) * isect.wt;
 }
 
-[numthreads(1000, 1, 1)]
+[numthreads(256, 1, 1)]
 void computeMain(uint3 threadIdx : SV_DispatchThreadID,)
 {
     uint seed = (threadIdx.x * threadIdx.x) * 30 + 3;
@@ -284,4 +285,15 @@ void computeMain(uint3 threadIdx : SV_DispatchThreadID,)
 // Expect: Approximately 0.0 in endpointDifferentialBuffer[2]
 // Expect: Approximately 0.0 in endpointDifferentialBuffer[3]
 //
-}
\ No newline at end of file
+}
+// CHECK: type: float
+// CHECK-NEXT: -0.{{9[5-9][0-9]}}000
+// CHECK-NEXT: 0.{{9[5-9][0-9]}}000
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 0.004000
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 0.000000
+// CHECK-NEXT: 0.000000
diff --git a/tests/autodiff/was/warped-sampling-1d.slang.expected.txt b/tests/autodiff/was/warped-sampling-1d.slang.expected.txt
deleted file mode 100644
index 84272b50a..000000000
--- a/tests/autodiff/was/warped-sampling-1d.slang.expected.txt
+++ /dev/null
@@ -1,11 +0,0 @@
-type: float
--0.954000
-0.950000
--0.000000
-0.004000
-0.000000
-0.000000
-0.000000
-0.000000
-0.000000
-0.000000
\ No newline at end of file