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>

2 files changed, 15 insertions, 6 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
+}