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-basic.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-basic.slang')
| -rw-r--r-- | tests/autodiff/autopybind-basic.slang | 25 |
1 files changed, 25 insertions, 0 deletions
diff --git a/tests/autodiff/autopybind-basic.slang b/tests/autodiff/autopybind-basic.slang new file mode 100644 index 000000000..5d4409474 --- /dev/null +++ b/tests/autodiff/autopybind-basic.slang @@ -0,0 +1,25 @@ +//TEST:SIMPLE(filecheck=CUDA): -target cuda -line-directive-mode none +//TEST:SIMPLE(filecheck=TORCH): -target torch -line-directive-mode none + +// CUDA: __global__ void __kernel__myKernel(TensorView inValues_[[#]], TensorView outValues_[[#]]) +[AutoPyBindCUDA] +[CudaKernel] +void myKernel(TensorView<float> inValues, TensorView<float> outValues) +{ + if (cudaThreadIdx().x > 0) + return; + outValues.store(cudaThreadIdx().x, sin(inValues.load(cudaThreadIdx().x))); +} + +// TORCH: {{^SLANG_PRELUDE_EXPORT$}} +// TORCH-NEXT: void __kernel__myKernel(TensorView {{[[: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:]_]+}}, 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: m.def("myKernel", &myKernel, "myKernel"); + +// TORCH: m.def("__funcinfo__myKernel", &__funcinfo__myKernel, "__funcinfo__myKernel");
\ No newline at end of file |