diff options
| author | Sai Praveen Bangaru <31557731+saipraveenb25@users.noreply.github.com> | 2023-09-19 18:51:24 -0400 |
|---|---|---|
| committer | GitHub <noreply@github.com> | 2023-09-19 18:51:24 -0400 |
| commit | 739c3a7b53dc6489065fcd5e9f0a04370c5f9c8f (patch) | |
| tree | 593c86cbc184476479c66554cc6784b454bdec66 /tests/autodiff/autopybind-struct.slang | |
| parent | 359fdc9d556b4c493c588c5b8f93df85933634f8 (diff) | |
Added `[AutoPyBindCUDA]` for automatic kernel binding + `[PyExport]` for exporting type information (#3209)
* Initial: add a DiffTensor impl
* Auto-binding and diff tensor implementations now work
* Refactored diff-tensor implementation + added py-export for struct types
* Cleanup
* Update slang-ir-pytorch-cpp-binding.cpp
* Updated test names
* Update autodiff-data-flow.slang.expected
* Add more versions of load/store & default generic args for DiffTensorView.
* Add diagnostic for default generic arg and more tests
* Add more `[AutoPyBind]` tests
Diffstat (limited to 'tests/autodiff/autopybind-struct.slang')
| -rw-r--r-- | tests/autodiff/autopybind-struct.slang | 35 |
1 files changed, 35 insertions, 0 deletions
diff --git a/tests/autodiff/autopybind-struct.slang b/tests/autodiff/autopybind-struct.slang new file mode 100644 index 000000000..d94218494 --- /dev/null +++ b/tests/autodiff/autopybind-struct.slang @@ -0,0 +1,35 @@ +//TEST:SIMPLE(filecheck=CUDA): -target cuda -line-directive-mode none +//TEST:SIMPLE(filecheck=TORCH): -target torch -line-directive-mode none + +[PyExport("Foo")] +struct Foo +{ + TensorView<float> inValues1; + TensorView<float> inValues2; +} + +// CUDA: __global__ void __kernel__myKernel(Foo_[[#]] {{[[:alnum:]_]+}}, TensorView {{[[:alnum:]_]+}}) +[AutoPyBindCUDA] +[CudaKernel] +void myKernel(Foo foo, TensorView<float> outValues) +{ + if (cudaThreadIdx().x > 0) + return; + outValues.store( + cudaThreadIdx().x, + sin(foo.inValues1.load(cudaThreadIdx().x)) * cos(foo.inValues2.load(cudaThreadIdx().x))); +} + +// TORCH: {{^SLANG_PRELUDE_EXPORT$}} +// TORCH-NEXT: __kernel__myKernel(Foo_[[#]] {{[[:alnum:]_]+}}, TensorView {{[[:alnum:]_]+}}); +// +// +// TORCH: {{^SLANG_PRELUDE_EXPORT$}} +// 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, torch::Tensor> {{[[:alnum:]_]+}}, torch::Tensor {{[[:alnum:]_]+}}) +// +// TORCH: {{^SLANG_PRELUDE_EXPORT$}} +// 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() +// +// TORCH: {{^SLANG_PRELUDE_EXPORT$}} +// TORCH-NEXT: std::tuple<std::tuple<const char*, const char*>, std::tuple<const char*, const char*>> __typeinfo__Foo() + |