diff options
Diffstat (limited to 'tests/autodiff')
| -rw-r--r-- | tests/autodiff/cuda-kernel-export-2.slang | 7 | ||||
| -rw-r--r-- | tests/autodiff/cuda-kernel-export.slang | 14 |
2 files changed, 15 insertions, 6 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 +} |
