yum-mirror/slang

Making it easier to work with shaders

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

Sai Praveen BangaruRemove `IRHLSLExportDecoration` and `IRKeepAliveDecoration` for non-CUDA/Torch targets (#4364)fba316f0e

master
branchesmasterall refs →
1.4 KiB55 linesraw
1//TEST:SIMPLE(filecheck=HLSL): -target hlsl -line-directive-mode none -entry computeMain -stage compute
2//TEST:SIMPLE(filecheck=CUDA): -target cuda -line-directive-mode none
3//TEST:SIMPLE(filecheck=TORCH): -target torch -line-directive-mode none
4
5//TEST_INPUT:ubuffer(data=[0 0 0 0 0], stride=4):out,name=outputBuffer
6RWStructuredBuffer<float> outputBuffer;
7
8typedef DifferentialPair<float> dpfloat;
9typedef float.Differential dfloat;
10
11[Differentiable]
12float func1(float x)
13{
14    return x * 4;
15}
16
17[AutoPyBindCUDA]
18[CUDAKernel]
19void torchMain(TensorView<float> v)
20{
21    v[0] = func1(v[0]);
22    v[1] = func1(v[1]);
23}
24
25// Shouldn't see torchMain (or its transformations) anywhere in the HLSL output
26// HLSL-NOT:torchMain
27// HLSL:func1
28// HLSL-NOT:torchMain
29// HLSL:computeMain
30// HLSL-NOT:torchMain
31
32[Differentiable]
33float func2(float a)
34{
35    return a;
36}
37
38[numthreads(1, 1, 1)]
39void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID)
40{
41    {
42        dpfloat dpa = dpfloat(2.0, 1.0);
43        dpfloat dpb = dpfloat(1.5, 1.0);
44
45        outputBuffer[0] = fwd_diff(func1)(dpa).d;                 // Expect: 1
46        outputBuffer[1] = fwd_diff(func2)(dpfloat(dpa.p, 0.0)).d;   // Expect: 0
47    }
48}
49
50// Ensure that the generated CUDA and Torch kernels do have torchMain & its transformations
51
52// TORCH:      {{^SLANG_PRELUDE_EXPORT$}}
53// TORCH-NEXT: void __kernel__torchMain(TensorView {{[[:alnum:]_]+}});
54
55// CUDA: __global__ void __kernel__torchMain(TensorView {{[[:alnum:]_]+}})