Implement FileCheck tests for several test commands (#2747)

* Add missing expected.txt for test

* Diagnostics -> StdWriters in render test

* Allow specifying several test prefixes to run

`slang-test -- tests/foo tests/bar`

* Squash warnings in some tests

* Enable gfx debug layer in gfx test util

Makes this issue present consistently: https://github.com/shader-slang/slang/issues/2766

* Allow DebugDevice to return interfaces instantiated by the debugged object

* Check that we actaully have a shader cache for shader cache tests

* Implement FileCheck tests for several test commands

- SIMPLE, SIMPLE_EX
- SIMPLE_LINE
- REFLECTION, CPU_REFLECTION
- CROSS_COMPILE

It does not currently support the render tests or the COMPARE_COMPUTE commands

It is invoked by adding `(filecheck=MY_FILECHECK_PREFIX)` to the test command, for example

TEST:CROSS_COMPILE(filecheck=SPIRV): -target spirv-assembly

* Move LLVM FileCheck interface to slang-llvm

* Neaten slang-test tests

* Refine handling of expected output in slang-test

* Add example FileCheck buffer test

* Add cuda-kernel-export tests

Which were waiting on FileCheck

* Bump vs project files

* Make createLLVMFileCheck_V1 return a void* rather than specifically an IFileCheck

* Remove use of CharSlice from filecheck interface

* Bump slang-llvm version

---------

Co-authored-by: jsmall-nvidia <jsmall@nvidia.com>

8 files changed, 47 insertions, 9 deletions

diff --git a/tests/autodiff/cuda-kernel-export-2.slang b/tests/autodiff/cuda-kernel-export-2.slang
index 9cbb4e881..93abddfa2 100644
--- a/tests/autodiff/cuda-kernel-export-2.slang
+++ b/tests/autodiff/cuda-kernel-export-2.slang
@@ -1,13 +1,13 @@
-//DISABLE_TEST:SIMPLE: -target cuda -line-directive-mode none
+//TEST:SIMPLE(filecheck=CUDA): -target cuda -line-directive-mode none
 
 // Verify that we can output a cuda device function with [CudaDeviceExport].
-// Disabled until we have FileCheck.
 
 
 //////////////////////////////////////////////////////////////////////////
 // Lambda GGX
 //////////////////////////////////////////////////////////////////////////
 
+// CUDA-DAG: __device__ float lambdaGGX(float alphaSqr_[[#]], float cosTheta_[[#]])
 [CudaDeviceExport]
 [BackwardDifferentiable]
 float lambdaGGX(const float alphaSqr, const float cosTheta)
@@ -19,6 +19,7 @@ float lambdaGGX(const float alphaSqr, const float cosTheta)
     return 0.5f * (sqrt(1.0f + alphaSqr * tanThetaSqr) - 1.0f);
 }
 
+// CUDA-DAG: __device__ void lambdaGGX_bwd(DiffPair_float_[[#]] * alphaSqr_[[#]], DiffPair_float_[[#]] * cosTheta_[[#]], float d_out_[[#]])
 [CudaDeviceExport]
 void lambdaGGX_bwd(inout DifferentialPair<float> alphaSqr, inout DifferentialPair<float> cosTheta, const float d_out)
 {
@@ -29,6 +30,7 @@ void lambdaGGX_bwd(inout DifferentialPair<float> alphaSqr, inout DifferentialPai
 // Masking Smith
 //////////////////////////////////////////////////////////////////////////
 
+// CUDA-DAG: __device__ float maskingSmithGGXCorrelated(float alphaSqr_[[#]], float cosThetaI_[[#]], float cosThetaO_[[#]])
 [CudaDeviceExport]
 [BackwardDifferentiable]
 float maskingSmithGGXCorrelated(const float alphaSqr, const float cosThetaI, const float cosThetaO)
@@ -38,6 +40,7 @@ float maskingSmithGGXCorrelated(const float alphaSqr, const float cosThetaI, con
     return 1.0f / (1.0f + lambdaI + lambdaO);
 }
 
+// CUDA-DAG: __device__ void maskingSmithGGXCorrelated_bwd(DiffPair_float_[[#]] * alphaSqr_[[#]], DiffPair_float_[[#]] * cosThetaI_[[#]], DiffPair_float_[[#]] * cosThetaO_[[#]], float d_out_[[#]])
 [CudaDeviceExport]
 void maskingSmithGGXCorrelated_bwd(inout DifferentialPair<float> alphaSqr,
                                    inout DifferentialPair<float> cosThetaI,
diff --git a/tests/autodiff/cuda-kernel-export.slang b/tests/autodiff/cuda-kernel-export.slang
index e16188abc..928133c94 100644
--- a/tests/autodiff/cuda-kernel-export.slang
+++ b/tests/autodiff/cuda-kernel-export.slang
@@ -1,7 +1,7 @@
-//DISABLE_TEST:SIMPLE: -target cuda -line-directive-mode none
+//TEST:SIMPLE(filecheck=CUDA): -target cuda -line-directive-mode none
+//TEST:SIMPLE(filecheck=TORCH): -target torch -line-directive-mode none
 
-// Verify that we can output a cuda device function with [CudaDeviceExport].
-// Disabled until we have FileCheck.
+// Verify that we can output a cuda device function with [CudaKernel].
 
 struct MySubType
 {
@@ -20,6 +20,7 @@ struct MyInput
     float normalVal;
 }
 
+// CUDA: __global__ void myKernel(TensorView inValues_[[#]], TensorView outValues_[[#]])
 [CudaKernel]
 void myKernel(TensorView<float> inValues, TensorView<float> outValues)
 {
@@ -28,6 +29,11 @@ void myKernel(TensorView<float> inValues, TensorView<float> outValues)
     outValues.store(cudaThreadIdx().x, sin(inValues.load(cudaThreadIdx().x)));
 }
 
+// TORCH:      {{^SLANG_PRELUDE_EXPORT$}}
+// TORCH-NEXT: void myKernel(TensorView {{[[:alnum:]_]+}}, TensorView {{[[:alnum:]_]+}});
+//
+// TORCH:      {{^SLANG_PRELUDE_EXPORT$}}
+// TORCH-NEXT: std::tuple<std::tuple<float, float>, std::tuple<std::tuple<std::tuple<torch::Tensor, torch::Tensor>>, std::tuple<std::tuple<torch::Tensor, torch::Tensor>>>> runCompute(std::tuple<torch::Tensor, float> input_[[#]])
 [TorchEntryPoint]
 public __extern_cpp MyType runCompute(MyInput input)
 {
@@ -44,4 +50,4 @@ public __extern_cpp MyType runCompute(MyInput input)
     rs.sub[1].array[0] = inValues;
     rs.sub[1].array[1] = outValues;
     return rs;
-}
\ No newline at end of file
+}
diff --git a/tests/bugs/gh-449.slang b/tests/bugs/gh-449.slang
index 9ce678b53..1b7a9f07d 100644
--- a/tests/bugs/gh-449.slang
+++ b/tests/bugs/gh-449.slang
@@ -1,5 +1,6 @@
 // gh-449.slang
 //TEST:SIMPLE:
+//TEST:SIMPLE(filecheck=CHECK):
 
 // Issue when dealing with binary operations that
 // mix scalars with vectors that have a different
@@ -16,10 +17,16 @@ void main()
     float2 a = float2(1, 2);
     uint   b = 3;
     foo(a + b);
+    // CHECK:      tests/bugs/gh-449.slang([[#@LINE-1]]): error 30019: expected an expression of type 'S', got 'vector<float,2>'
+    // CHECK-NEXT:     foo(a + b);
+    // CHECK-NEXT:           ^
 
     // This used to get confused, with the `f` getting converted
     // to a `uint` before the addition.
     uint2 u = uint2(1, 2);
     float f = 3.0;
     foo(u + f);
+    // CHECK:      tests/bugs/gh-449.slang([[#@LINE-1]]): error 30019: expected an expression of type 'S', got 'vector<float,2>'
+    // CHECK-NEXT:     foo(u + f);
+    // CHECK-NEXT:           ^
 }
diff --git a/tests/bugs/gh-449.slang.expected b/tests/bugs/gh-449.slang.expected
index 8ca87a406..4d01e7fe3 100644
--- a/tests/bugs/gh-449.slang.expected
+++ b/tests/bugs/gh-449.slang.expected
@@ -1,9 +1,9 @@
 result code = -1
 standard error = {
-tests/bugs/gh-449.slang(18): error 30019: expected an expression of type 'S', got 'vector<float,2>'
+tests/bugs/gh-449.slang(19): error 30019: expected an expression of type 'S', got 'vector<float,2>'
     foo(a + b);
           ^
-tests/bugs/gh-449.slang(24): error 30019: expected an expression of type 'S', got 'vector<float,2>'
+tests/bugs/gh-449.slang(28): error 30019: expected an expression of type 'S', got 'vector<float,2>'
     foo(u + f);
           ^
 }
diff --git a/tests/bugs/gh-841.slang b/tests/bugs/gh-841.slang
index 44d8348e5..8dc687e77 100644
--- a/tests/bugs/gh-841.slang
+++ b/tests/bugs/gh-841.slang
@@ -1,12 +1,26 @@
 // gh-841.slang
 
 //TEST:CROSS_COMPILE: -profile ps_5_0 -entry main -target spirv-assembly
+//TEST:CROSS_COMPILE(filecheck=SPV): -profile ps_5_0 -entry main -target spirv-assembly
+//TEST:CROSS_COMPILE(filecheck=GLSL): -profile ps_5_0 -entry main -target glsl
 
 // GitHub issue #841: failing to emit `flat` modifier in output GLSL when required
 
 struct RasterVertex
 {
+  // location 0
 	float4 c : COLOR;
+
+  // Make sure that the input value in location 1 is decorated as Flat
+  // SPV-DAG: [[#VAL:]]{{.*}}:{{.*}} Variable Input
+  // SPV-DAG: Decorate [[#VAL]]{{.*}} Location 1
+  // SPV-DAG: Decorate [[#VAL]]{{.*}} Flat
+  //
+  // Likewise for GLSL
+  // GLSL:      flat layout(location = 1)
+  // GLSL-NEXT: in uint {{[[:alnum:]_]+}};
+  //
+  // location 1
 	uint u : FLAGS;
 };
 
diff --git a/tests/compute/simple.slang b/tests/compute/simple.slang
index c62f88d5e..110a590a1 100644
--- a/tests/compute/simple.slang
+++ b/tests/compute/simple.slang
@@ -1,5 +1,11 @@
 //TEST(smoke,compute):COMPARE_COMPUTE: -shaderobj
 //TEST(smoke,compute):COMPARE_COMPUTE:-cpu -shaderobj
+//TEST(smoke,compute):COMPARE_COMPUTE(filecheck-buffer=CHECK):-cpu -shaderobj
+
+// CHECK:      0
+// CHECK-NEXT: 3F800000
+// CHECK-NEXT: 40000000
+// CHECK-NEXT: 40400000
 
 // This is a basic test for Slang compute shader.
 
diff --git a/tests/diagnostics/local-line.slang b/tests/diagnostics/local-line.slang
index 973c234c0..532307a8f 100644
--- a/tests/diagnostics/local-line.slang
+++ b/tests/diagnostics/local-line.slang
@@ -3,6 +3,7 @@
 //TEST:SIMPLE_LINE:-entry computeMain -target dxbc  -stage compute
 //TEST:SIMPLE_LINE:-entry computeMain -target dll -stage compute
 //TEST:SIMPLE_LINE:-entry computeMain -target ptx -stage compute
+//TEST:SIMPLE_LINE(filecheck=CHECK):-entry computeMain -target spirv -stage compute
 
 //TEST_INPUT:ubuffer(data=[0 0 0 0], stride=4):out,name outputBuffer
 RWStructuredBuffer<int> outputBuffer;
@@ -37,6 +38,7 @@ void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID)
         else
         {
             a = e; b = c + c; d += d + __SyntaxError(); e = doThing(e, int(dispatchThreadID.x));
+            // CHECK: [[#@LINE-1]]
             
         }
     }
diff --git a/tests/diagnostics/local-line.slang.expected b/tests/diagnostics/local-line.slang.expected
index a2720097d..425151f3a 100644
--- a/tests/diagnostics/local-line.slang.expected
+++ b/tests/diagnostics/local-line.slang.expected
@@ -1 +1 @@
-39
+40