From 0a1826034d17304422d1f3c0bde66c2c5da0420c Mon Sep 17 00:00:00 2001
From: Corey Lowman <clowman1993@gmail.com>
Date: Wed, 5 Jul 2023 20:20:41 +0000
Subject: [PATCH] Fixing compilation errors

---
 src/tensor_ops/convtrans2d/convtrans2d.cu | 4 ++--
 src/tensor_ops/convtrans2d/cpu_kernel.rs  | 1 -
 src/tensor_ops/convtrans2d/cuda_kernel.rs | 5 ++---
 3 files changed, 4 insertions(+), 6 deletions(-)

diff --git a/src/tensor_ops/convtrans2d/convtrans2d.cu b/src/tensor_ops/convtrans2d/convtrans2d.cu
index be5b21f83..3027e8d8d 100644
--- a/src/tensor_ops/convtrans2d/convtrans2d.cu
+++ b/src/tensor_ops/convtrans2d/convtrans2d.cu
@@ -129,11 +129,11 @@ __device__ void transpose_filters(
     filters_tr += k1 * op.kernel;
     filters_tr += cg * (op.kernel * op.kernel);
     filters_tr += og * (c_per_g * op.kernel * op.kernel);
-    filters_tr += g * (o_per_g * * c_per_g * op.kernel * op.kernel);
+    filters_tr += g * (o_per_g * c_per_g * op.kernel * op.kernel);
     *filters_tr = filters[i_no];
 }
 
-#define CONV_OP(TYPENAME, UNFOLD_INPUT, UNFOLD_OUTPUT, TR_FILTERS, SUM_TR_FILTERS) \
+#define CONV_OP(TYPENAME, UNFOLD_INPUT, UNFOLD_OUTPUT, TR_FILTERS) \
 extern "C" __global__ void UNFOLD_INPUT( \
     const Conv2DOp op, \
     const TYPENAME *image, \
diff --git a/src/tensor_ops/convtrans2d/cpu_kernel.rs b/src/tensor_ops/convtrans2d/cpu_kernel.rs
index 5c454c6fe..f61943f84 100644
--- a/src/tensor_ops/convtrans2d/cpu_kernel.rs
+++ b/src/tensor_ops/convtrans2d/cpu_kernel.rs
@@ -254,7 +254,6 @@ where
         let lhs = lhs.data.as_ref();
 
         let rhs = rhs.data.as_ref();
-        let grad_rhs = grad_rhs.data.as_mut();
         for i_batch in 0..op.batch {
             self.convtrans2d_backward(
                 &op,
diff --git a/src/tensor_ops/convtrans2d/cuda_kernel.rs b/src/tensor_ops/convtrans2d/cuda_kernel.rs
index b4bb947bd..827567220 100644
--- a/src/tensor_ops/convtrans2d/cuda_kernel.rs
+++ b/src/tensor_ops/convtrans2d/cuda_kernel.rs
@@ -100,7 +100,7 @@ where
             // generate patches for matmul
             let unfold_fn = self.dev.get_func(Self::MOD, Self::FNS[0]).unwrap();
             let cfg = launch_cfg::<128>((op.batch * op.chan_in * op.h_out * op.w_out) as u32);
-            unsafe { unfold_fn.launch(cfg, (op, lhs.data.as_ref(), &img_strides, &mut patches)) }?;
+            unfold_fn.launch(cfg, (op, lhs.data.as_ref(), &img_strides, &mut patches))?;
 
             // prepare filters for backward operations by
             // swapping dims 0 and 1 and adding a batch dimension
@@ -163,7 +163,6 @@ where
         }
 
         let rhs = rhs.data.as_ref();
-        let grad_rhs = grad_rhs.data.as_mut();
 
         unsafe {
             self.par_stream.wait_for_default()?;
@@ -238,7 +237,7 @@ where
                         &patches.slice(i_batch * op.groups * k * n..),
                         [k * n, 1, k],
                         Default::default(),
-                        grad_rhs.slice_mut(i_batch * op.groups * m * n..),
+                        &mut grad_rhs.slice_mut(i_batch * op.groups * m * n..),
                         [m * n, n, 1],
                     )
                     .unwrap();