From d6dd38f5210fedfd96ff088ef9b8a2424c79c4f0 Mon Sep 17 00:00:00 2001 From: Ellie Hermaszewska Date: Tue, 11 Apr 2023 23:28:58 +0800 Subject: 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 --- tests/autodiff/cuda-kernel-export-2.slang | 7 +++++-- tests/autodiff/cuda-kernel-export.slang | 14 ++++++++++---- 2 files changed, 15 insertions(+), 6 deletions(-) (limited to 'tests/autodiff') 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 alphaSqr, inout DifferentialPair cosTheta, const float d_out) { @@ -29,6 +30,7 @@ void lambdaGGX_bwd(inout DifferentialPair 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 alphaSqr, inout DifferentialPair 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 inValues, TensorView outValues) { @@ -28,6 +29,11 @@ void myKernel(TensorView inValues, TensorView 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>, std::tuple>>> runCompute(std::tuple 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 +} -- cgit v1.2.3