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>

1 files changed, 5 insertions, 2 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,