Backward convolution implemented.

Author: vbharadwaj-bk

openequivariance_extjax/src/libjax_tp_jit.cpp

@@ -471,6 +471,64 @@ ffi::Error conv_forward_impl(
     return ffi::Error::Success();
 }
 
+ffi::Error conv_backward_impl(
+        ffi::AnyBuffer L1_in,
+        ffi::AnyBuffer L2_in,
+        ffi::AnyBuffer W,
+        ffi::AnyBuffer L3_grad,
+        ffi::Result<ffi::AnyBuffer> L1_grad,
+        ffi::Result<ffi::AnyBuffer> L2_grad,
+        ffi::Result<ffi::AnyBuffer> W_grad, 
+        ffi::AnyBuffer rows,
+        ffi::AnyBuffer cols,
+        ffi::AnyBuffer workspace,
+        ffi::AnyBuffer transpose_perm,
+        cudaStream_t stream, 
+        std::string_view kernel, ffi::Dictionary forward_config, ffi::Dictionary backward_config, ffi::Dictionary double_backward_config, ffi::Dictionary kernel_prop,
+        int64_t hash) {
+    
+    auto [jit_kernel, k] = compile_conv_with_caching(
+        kernel, forward_config, backward_config, double_backward_config, kernel_prop, hash, true);
+    const int64_t nnz = rows.dimensions()[0];
+    const int64_t node_count = L1_in.dimensions()[0];
+    check_tensor(L1_in, {node_count, k.L1_dim}, k.irrep_dtype, "L1_in");
+    check_tensor(L2_in, {nnz, k.L2_dim}, k.irrep_dtype, "L2_in");
+    check_tensor(L3_grad, {node_count, k.L3_dim}, k.irrep_dtype, "L3_grad");
+    check_tensor(workspace, {k.workspace_size}, k.workspace_dtype, "workspace");
+    check_tensor(rows, {nnz}, k.idx_dtype, "rows");
+    check_tensor(cols, {nnz}, k.idx_dtype, "cols");
+
+    if (k.deterministic)
+        check_tensor(transpose_perm, {nnz}, k.idx_dtype, "transpose perm");
+    
+    if (k.shared_weights) {
+        check_tensor(W, {k.weight_numel}, k.weight_dtype, "W");
+        check_tensor(*W_grad, {k.weight_numel}, k.weight_dtype, "W_grad");
+    }
+    else {
+        check_tensor(W, {nnz, k.weight_numel}, k.weight_dtype, "W");
+        check_tensor(*W_grad, {nnz, k.weight_numel}, k.weight_dtype, "W_grad");
+    }
+    if(k.shared_weights)
+        zero_buffer(*W_grad);
+
+    jit_kernel->backward(
+            data_ptr(L1_in),
+            data_ptr(L1_grad),
+            data_ptr(L2_in),
+            data_ptr(L2_grad),
+            data_ptr(W),
+            data_ptr(W_grad),
+            data_ptr(L3_grad),
+            data_ptr(rows),
+            data_ptr(cols),
+            nnz, node_count,
+            data_ptr(workspace),
+            data_ptr(transpose_perm),
+            stream);
+    return ffi::Error::Success();
+}
+
 XLA_FFI_DEFINE_HANDLER_SYMBOL(
     conv_forward, conv_forward_impl,
     ffi::Ffi::Bind()
@@ -487,6 +545,25 @@ XLA_FFI_DEFINE_HANDLER_SYMBOL(
         .Attr<int64_t>("hash"),
         {xla::ffi::Traits::kCmdBufferCompatible});
 
+XLA_FFI_DEFINE_HANDLER_SYMBOL(
+    conv_backward, conv_backward_impl,
+    ffi::Ffi::Bind()
+        .Arg<ffi::AnyBuffer>()
+        .Arg<ffi::AnyBuffer>()
+        .Arg<ffi::AnyBuffer>()
+        .Arg<ffi::AnyBuffer>()
+        .Ret<ffi::AnyBuffer>()
+        .Ret<ffi::AnyBuffer>()
+        .Ret<ffi::AnyBuffer>()
+        .Arg<ffi::AnyBuffer>()
+        .Arg<ffi::AnyBuffer>()
+        .Arg<ffi::AnyBuffer>()
+        .Arg<ffi::AnyBuffer>()
+        .Ctx<ffi::PlatformStream<cudaStream_t>>()
+        .Attr<std::string_view>("kernel").Attr<ffi::Dictionary>("forward_config").Attr<ffi::Dictionary>("backward_config").Attr<ffi::Dictionary>("double_backward_config").Attr<ffi::Dictionary>("kernel_prop")
+        .Attr<int64_t>("hash"),
+        {xla::ffi::Traits::kCmdBufferCompatible});
+
 NB_MODULE(openequivariance_extjax, m) {
     m.def("registrations", []() {
         nb::dict registrations;