yum-mirror/slang

Making it easier to work with shaders

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

Sai Praveen BangaruFix lowering of associated types in generic interfaces (#6600)78517dc39

master
branchesmasterall refs →
2.9 KiB47 linesraw
1//TEST:SIMPLE(filecheck=CUDA): -target cuda -line-directive-mode none
2//TEST:SIMPLE(filecheck=TORCH): -target torch -line-directive-mode none
3
4// CUDA: __device__ void s_primal_ctx_myKernel_0(
5// CUDA:         printf("%f\n",
6// CUDA: __global__ void __kernel__myKernel_bwd_diff(DiffTensorView_[[#]] {{[[:alnum:]_]+}}, DiffTensorView_[[#]] {{[[:alnum:]_]+}})
7// CUDA: __global__ void __kernel__myKernel_fwd_diff(DiffTensorView_[[#]] {{[[:alnum:]_]+}}, DiffTensorView_[[#]] {{[[:alnum:]_]+}})
8// CUDA: __global__ void __kernel__myKernel(DiffTensorView_[[#]] {{[[:alnum:]_]+}}, DiffTensorView_[[#]] {{[[:alnum:]_]+}})
9
10[AutoPyBindCUDA]
11[Differentiable]
12[CudaKernel]
13void myKernel(DiffTensorView inValues, DiffTensorView outValues)
14{
15    if (cudaThreadIdx().x > 0)
16        return;
17    printf("%f\n", inValues[cudaThreadIdx().x]);
18    outValues[cudaThreadIdx().x] = sin(inValues[cudaThreadIdx().x]);
19}
20
21// TORCH:      {{^SLANG_PRELUDE_EXPORT$}}
22// TORCH-NEXT: void __kernel__myKernel_bwd_diff(DiffTensorView_[[#]] {{[[:alnum:]_]+}}, DiffTensorView_[[#]] {{[[:alnum:]_]+}})
23//
24// TORCH:      {{^SLANG_PRELUDE_EXPORT$}}
25// TORCH-NEXT: void __kernel__myKernel_fwd_diff(DiffTensorView_[[#]] {{[[:alnum:]_]+}}, DiffTensorView_[[#]] {{[[:alnum:]_]+}})
26//
27// TORCH:      {{^SLANG_PRELUDE_EXPORT$}}
28// TORCH-NEXT: void __kernel__myKernel(DiffTensorView_[[#]] {{[[:alnum:]_]+}}, DiffTensorView_[[#]] {{[[:alnum:]_]+}})
29//
30// TORCH:      {{^SLANG_PRELUDE_EXPORT$}}
31// TORCH-NEXT: void myKernel(std::tuple<uint32_t, uint32_t, uint32_t> {{[[:alnum:]_]+}}, std::tuple<uint32_t, uint32_t, uint32_t> {{[[:alnum:]_]+}}, std::tuple<torch::Tensor, std::tuple<torch::Tensor>> {{[[:alnum:]_]+}}, std::tuple<torch::Tensor, std::tuple<torch::Tensor>> {{[[:alnum:]_]+}})
32//
33// TORCH:      {{^SLANG_PRELUDE_EXPORT$}}
34// TORCH-NEXT: std::tuple<std::tuple<const char*, const char*, const char*, const char*>, std::tuple<const char*, const char*>, const char*, const char*> __funcinfo__myKernel()
35//
36// TORCH:      {{^SLANG_PRELUDE_EXPORT$}}
37// TORCH-NEXT: void myKernel_fwd_diff(std::tuple<uint32_t, uint32_t, uint32_t> {{[[:alnum:]_]+}}, std::tuple<uint32_t, uint32_t, uint32_t> {{[[:alnum:]_]+}}, std::tuple<torch::Tensor, std::tuple<torch::Tensor>> {{[[:alnum:]_]+}}, std::tuple<torch::Tensor, std::tuple<torch::Tensor>> {{[[:alnum:]_]+}})
38//
39// TORCH:      {{^SLANG_PRELUDE_EXPORT$}}
40// TORCH-NEXT: void myKernel_bwd_diff(std::tuple<uint32_t, uint32_t, uint32_t> {{[[:alnum:]_]+}}, std::tuple<uint32_t, uint32_t, uint32_t> {{[[:alnum:]_]+}}, std::tuple<torch::Tensor, std::tuple<torch::Tensor>> {{[[:alnum:]_]+}}, std::tuple<torch::Tensor, std::tuple<torch::Tensor>> {{[[:alnum:]_]+}})
41//
42// TORCH:     {{^SLANG_PRELUDE_EXPORT$}}
43// TORCH-NEXT: std::tuple<std::tuple<const char*, const char*>, std::tuple<const char*, const char*>> __typeinfo__DiffTensorView()
44//
45// TORCH:     {{^SLANG_PRELUDE_EXPORT$}}
46// TORCH-NEXT: std::tuple<std::tuple<const char*>, std::tuple<const char*>> __typeinfo__AtomicAdd()
47//