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 ++++++++++---- tests/bugs/gh-449.slang | 7 +++++++ tests/bugs/gh-449.slang.expected | 4 ++-- tests/bugs/gh-841.slang | 14 ++++++++++++++ tests/compute/simple.slang | 6 ++++++ tests/diagnostics/local-line.slang | 2 ++ tests/diagnostics/local-line.slang.expected | 2 +- 8 files changed, 47 insertions(+), 9 deletions(-) (limited to 'tests') 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 +} 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' + // 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' + // 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' +tests/bugs/gh-449.slang(19): error 30019: expected an expression of type 'S', got 'vector' foo(a + b); ^ -tests/bugs/gh-449.slang(24): error 30019: expected an expression of type 'S', got 'vector' +tests/bugs/gh-449.slang(28): error 30019: expected an expression of type 'S', got 'vector' 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 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 -- cgit v1.2.3