diff options
Diffstat (limited to 'tests')
| -rw-r--r-- | tests/autodiff/cuda-kernel-export-2.slang | 7 | ||||
| -rw-r--r-- | tests/autodiff/cuda-kernel-export.slang | 14 | ||||
| -rw-r--r-- | tests/bugs/gh-449.slang | 7 | ||||
| -rw-r--r-- | tests/bugs/gh-449.slang.expected | 4 | ||||
| -rw-r--r-- | tests/bugs/gh-841.slang | 14 | ||||
| -rw-r--r-- | tests/compute/simple.slang | 6 | ||||
| -rw-r--r-- | tests/diagnostics/local-line.slang | 2 | ||||
| -rw-r--r-- | tests/diagnostics/local-line.slang.expected | 2 |
8 files changed, 47 insertions, 9 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 +} 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<float,2>' + // 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<float,2>' + // 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<float,2>' +tests/bugs/gh-449.slang(19): error 30019: expected an expression of type 'S', got 'vector<float,2>' foo(a + b); ^ -tests/bugs/gh-449.slang(24): error 30019: expected an expression of type 'S', got 'vector<float,2>' +tests/bugs/gh-449.slang(28): error 30019: expected an expression of type 'S', got 'vector<float,2>' 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<int> 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 |
