From 739c3a7b53dc6489065fcd5e9f0a04370c5f9c8f Mon Sep 17 00:00:00 2001 From: Sai Praveen Bangaru <31557731+saipraveenb25@users.noreply.github.com> Date: Tue, 19 Sep 2023 18:51:24 -0400 Subject: 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 --- tests/autodiff/autopybind-basic.slang | 25 +++++++++++++++++++++++++ 1 file changed, 25 insertions(+) create mode 100644 tests/autodiff/autopybind-basic.slang (limited to 'tests/autodiff/autopybind-basic.slang') 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 inValues, TensorView 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 {{[[:alnum:]_]+}}, std::tuple {{[[:alnum:]_]+}}, torch::Tensor {{[[:alnum:]_]+}}, torch::Tensor {{[[:alnum:]_]+}}) + +// TORCH: {{^SLANG_PRELUDE_EXPORT$}} +// TORCH-NEXT: std::tuple, std::tuple, 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 -- cgit v1.2.3