yum-mirror/slang

Making it easier to work with shaders

git clone https://git.yummers.dev/yum-mirror/slang

Ellie HermaszewskaImplement FileCheck tests for several test commands (#2747)d6dd38f52

master
branchesmasterall refs →
2.3 KiB51 linesraw
1//TEST:SIMPLE(filecheck=CUDA): -target cuda -line-directive-mode none
2
3// Verify that we can output a cuda device function with [CudaDeviceExport].
4
5
6//////////////////////////////////////////////////////////////////////////
7// Lambda GGX
8//////////////////////////////////////////////////////////////////////////
9
10// CUDA-DAG: __device__ float lambdaGGX(float alphaSqr_[[#]], float cosTheta_[[#]])
11[CudaDeviceExport]
12[BackwardDifferentiable]
13float lambdaGGX(const float alphaSqr, const float cosTheta)
14{
15    const float SPECULAR_EPSILON = 1e-4f;
16    float _cosTheta = clamp(cosTheta, SPECULAR_EPSILON, 1.0f - SPECULAR_EPSILON);
17    float cosThetaSqr = _cosTheta * _cosTheta;
18    float tanThetaSqr = (1.0 - cosThetaSqr) / cosThetaSqr;
19    return 0.5f * (sqrt(1.0f + alphaSqr * tanThetaSqr) - 1.0f);
20}
21
22// CUDA-DAG: __device__ void lambdaGGX_bwd(DiffPair_float_[[#]] * alphaSqr_[[#]], DiffPair_float_[[#]] * cosTheta_[[#]], float d_out_[[#]])
23[CudaDeviceExport]
24void lambdaGGX_bwd(inout DifferentialPair<float> alphaSqr, inout DifferentialPair<float> cosTheta, const float d_out)
25{
26    __bwd_diff(lambdaGGX)(alphaSqr, cosTheta, d_out);
27}
28
29//////////////////////////////////////////////////////////////////////////
30// Masking Smith
31//////////////////////////////////////////////////////////////////////////
32
33// CUDA-DAG: __device__ float maskingSmithGGXCorrelated(float alphaSqr_[[#]], float cosThetaI_[[#]], float cosThetaO_[[#]])
34[CudaDeviceExport]
35[BackwardDifferentiable]
36float maskingSmithGGXCorrelated(const float alphaSqr, const float cosThetaI, const float cosThetaO)
37{
38    float lambdaI = lambdaGGX(alphaSqr, cosThetaI);
39    float lambdaO = lambdaGGX(alphaSqr, cosThetaO);
40    return 1.0f / (1.0f + lambdaI + lambdaO);
41}
42
43// CUDA-DAG: __device__ void maskingSmithGGXCorrelated_bwd(DiffPair_float_[[#]] * alphaSqr_[[#]], DiffPair_float_[[#]] * cosThetaI_[[#]], DiffPair_float_[[#]] * cosThetaO_[[#]], float d_out_[[#]])
44[CudaDeviceExport]
45void maskingSmithGGXCorrelated_bwd(inout DifferentialPair<float> alphaSqr,
46                                   inout DifferentialPair<float> cosThetaI,
47                                   inout DifferentialPair<float> cosThetaO,
48                                   const float d_out)
49{
50    __bwd_diff(maskingSmithGGXCorrelated)(alphaSqr, cosThetaI, cosThetaO, d_out);
51}