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 +++++-- 1 file changed, 5 insertions(+), 2 deletions(-) (limited to 'tests/autodiff/cuda-kernel-export-2.slang') 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, -- cgit v1.2.3