Repository: coreylowman/dfdx (language: rust)

Issue:
Add workspace for Cuda & Cpu devices
This would primarily be used by convolutions & anything else that needs to allocate things for forward/backward pass.

The API might look like:
```rust
// Cpu
pub(crate) fn workspace<E>(&self, num: usize) -> Option<&mut [E]>;

// Cuda
pub(crate) fn workspace<E>(&self, num: usize) -> Option<CudaViewMut<E>>;
```

**Blocking** questions:
1. How can we do this while maintaining `&self`?
2. How could we do this in a threadsafe manner? Both Cpu/Cuda can be shared across threads, but it wouldn't be safe for multiple threads to access the workspace at a time. Surround in a Mutex?
3. How do we have a workspace with just a number of bytes, but not impose a dtype on them? I.e. we don't know ahead of time when we are constructing the device what the type of the workspace would be. So we can't do something like `workspace: Mutex<Vec<E>>`.





Discussion / hints:
What would the workspace be for? What things would need to be allocated for convolutions?
For conv2d, both forward & backward it has to unfold the images into patches to do the matmuls.

The forward allocation is here https://github.com/coreylowman/dfdx/blob/main/src/tensor_ops/conv2d/cuda_kernel.rs#L66. The idea behind workspace is to not have to re-allocate this temporary buffer each time, and instead re-use a shared one.

This could also be used for dropout which has to allocate a buffer to hold the noise https://github.com/coreylowman/dfdx/blob/main/src/tensor_ops/dropout/cuda_kernel.rs#L44.


Fix patch:
diff --git a/Cargo.toml b/Cargo.toml
index 3a4d832df..26920ee6a 100644
--- a/Cargo.toml
+++ b/Cargo.toml
@@ -32,7 +32,7 @@ matrixmultiply = { version = "0.3.2", default-features = false, optional = true
 zip = { version = "0.6.2", default-features = false, optional = true }
 cblas-sys = { version = "0.1.4", default-features = false, optional = true }
 libc = { version = "0.2", default-features = false, optional = true }
-cudarc = { version = "0.9.0", default-features = false, optional = true, features = ["driver", "cublas"] }
+cudarc = { version = "0.9.2", default-features = false, optional = true, features = ["driver", "cublas"] }
 num-traits = { version = "0.2.15", default-features = false }
 safetensors = { version = "0.3", default-features = false, optional = true }
 memmap2 = { version = "0.5", default-features = false, optional = true }
diff --git a/src/tensor/cuda/device.rs b/src/tensor/cuda/device.rs
index a3641e077..0fd35c398 100644
--- a/src/tensor/cuda/device.rs
+++ b/src/tensor/cuda/device.rs
@@ -6,7 +6,12 @@ use cudarc::{
     cublas::{result::CublasError, CudaBlas},
     driver::{CudaDevice, CudaSlice, CudaStream, DeviceSlice, DriverError},
 };
-use std::{sync::Arc, vec::Vec};
+
+use std::sync::MutexGuard;
+use std::{
+    sync::{Arc, Mutex},
+    vec::Vec,
+};
 
 /// A Cuda device that enables constructing tensors on GPUs
 /// & running GPU kernels.
@@ -17,6 +22,7 @@ pub struct Cuda {
     pub(crate) blas: Arc<CudaBlas>,
     /// A second stream for kernels to optionally execute on.
     pub(crate) par_stream: Arc<CudaStream>,
+    pub(crate) workspace: Arc<Mutex<CudaSlice<u8>>>,
 }
 
 #[derive(Debug)]
@@ -67,11 +73,13 @@ impl Cuda {
         let dev = CudaDevice::new(ordinal)?;
         let blas = Arc::new(CudaBlas::new(dev.clone())?);
         let par_stream = Arc::new(dev.fork_default_stream()?);
+        let workspace = Arc::new(Mutex::new(dev.alloc_zeros::<u8>(0)?));
         Ok(Self {
             cpu,
             dev,
             blas,
             par_stream,
+            workspace,
         })
     }
 
@@ -90,6 +98,24 @@ impl Cuda {
     }
 }
 
+impl Cuda {
+    pub(crate) unsafe fn get_workspace<E>(
+        &self,
+        len: usize,
+    ) -> Result<MutexGuard<CudaSlice<u8>>, CudaError> {
+        let num_bytes_required = len * std::mem::size_of::<E>();
+        let mut workspace = self.workspace.as_ref().lock().unwrap();
+
+        // re-allocate a larger workspace
+        if workspace.num_bytes() < num_bytes_required {
+            // we are about to memset this to zero, so this is still okay
+            *workspace = unsafe { self.dev.alloc::<u8>(num_bytes_required) }?;
+        }
+
+        Ok(workspace)
+    }
+}
+
 impl std::fmt::Display for CudaError {
     fn fmt(&self, f: &mut core::fmt::Formatter<'_>) -> core::fmt::Result {
         write!(f, "{self:?}")
diff --git a/src/tensor_ops/conv2d/conv2d.cu b/src/tensor_ops/conv2d/conv2d.cu
index 1c0b4ea07..d2d215d66 100644
--- a/src/tensor_ops/conv2d/conv2d.cu
+++ b/src/tensor_ops/conv2d/conv2d.cu
@@ -19,8 +19,8 @@ __device__ void unfold_input_into_patches(
     T *patches // 6d (Batch, Channels, KernelSize, KernelSize, HeightOut, WidthOut)
 ) {
     unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
-    const auto patches_numel = op.batch * op.chan_in * op.kernel * op.kernel * op.h_out * op.w_out;
-    if (i >= patches_numel) {
+    size_t item_numel = op.batch * op.chan_in * op.kernel * op.kernel * op.h_out * op.w_out;
+    if (i >= item_numel) {
         return;
     }
 
@@ -37,28 +37,18 @@ __device__ void unfold_input_into_patches(
     const size_t c = idx % op.chan_in;
     idx /= op.chan_in;
     const size_t b = idx % op.batch;
-    idx /= op.batch;
 
     const size_t y_plus_p = oh * op.stride + k1;
-    if (y_plus_p < op.padding) {
-        return;
-    }
     const size_t y = y_plus_p - op.padding;
-    if (y >= op.h_in) {
-        return;
-    }
-
     const size_t x_plus_p = ow * op.stride + k2;
-    if (x_plus_p < op.padding) {
-        return;
-    }
     const size_t x = x_plus_p - op.padding;
-    if (x >= op.w_in) {
-        return;
-    }
 
-    const size_t i_image = b * strides[0] + c * strides[1] + y * strides[2] + x * strides[3];
-    patches[i] = image[i_image];
+    if (y >= op.h_in || x >= op.w_in) {
+        patches[i] = 0.0;
+    } else {
+        const size_t i_image = b * strides[0] + c * strides[1] + y * strides[2] + x * strides[3];
+        patches[i] = image[i_image];
+    }
 }
 
 template<typename T>
@@ -67,9 +57,9 @@ __device__ void unfold_output_into_patches(
     const T *image_out, // 4d (Batch, ChanOut, HeightOut, WidthOut)
     T *patches // 6d (Batch, ChanOut, KernelSize, KernelSize, HeightIn, WidthIn)
 ) {
-    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
-    const auto patches_numel = op.batch * op.chan_out * op.kernel * op.kernel * op.h_in * op.w_in;
-    if (i >= patches_numel) {
+    const unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
+    const size_t item_numel = op.chan_out * op.kernel * op.kernel * op.h_in * op.w_in;
+    if (i >= item_numel) {
         return;
     }
 
@@ -83,46 +73,36 @@ __device__ void unfold_output_into_patches(
     const size_t k1 = idx % op.kernel;
     idx /= op.kernel;
     const size_t o = idx % op.chan_out;
-    idx /= op.chan_out;
-    const size_t b = idx % op.batch;
-    idx /= op.batch;
 
-    size_t oh = y + op.padding;
-    if (oh < k1) {
-        return;
-    }
-    oh -= k1;
-    if (oh % op.stride != 0) {
-        return;
-    }
-    oh /= op.stride;
-    if (oh >= op.h_out) {
-        return;
-    }
-    
-    size_t ow = x + op.padding;
-    if (ow < k2) {
-        return;
-    }
-    ow -= k2;
-    if (ow % op.stride != 0) {
-        return;
-    }
-    ow /= op.stride;
-    if (ow >= op.w_out) {
+    const size_t oh_ks = y + op.padding;
+    const size_t oh_s = oh_ks - k1;
+    const size_t oh = oh_s / op.stride;
+    const size_t ow_ks = x + op.padding;
+    const size_t ow_s = ow_ks - k2;
+    const size_t ow = ow_s / op.stride;
+
+    if (
+        (oh_ks < k1 || oh_s % op.stride != 0 || oh >= op.h_out)
+        || (ow_ks < k2 || ow_s % op.stride != 0 || ow >= op.w_out)
+    ) {
+        for (auto b = 0; b < op.batch; b++) {
+            patches[b * item_numel + i] = 0.0;
+        }
         return;
     }
 
-    size_t image_i = b * (op.chan_out * op.h_out * op.w_out) + o * (op.h_out * op.w_out) + oh * (op.w_out)  + ow;
-    patches[i] = image_out[image_i];
+    for (auto b = 0; b < op.batch; b++) {
+        size_t image_i = b * (op.chan_out * op.h_out * op.w_out) + o * (op.h_out * op.w_out) + oh * (op.w_out)  + ow;
+        patches[b * item_numel + i] = image_out[image_i];
+    }
 }
 
 template<typename T>
-__device__ void transpose_and_broadcast_filters(
+__device__ void transpose_filters(
     const Conv2DOp op,
     const T *filters, // 4d (ChanOut, ChanIn, KernelSize, KernelSize)
     const size_t *strides, // 4d filters strides
-    T *filters_tr // 5d (Batch, ChanIn, ChanOut, KernelSize, KernelSize)
+    T *filters_tr // 4d (ChanIn, ChanOut, KernelSize, KernelSize)
 ) {
     unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
     auto numel = op.chan_in * op.chan_out * op.kernel * op.kernel;
@@ -138,15 +118,11 @@ __device__ void transpose_and_broadcast_filters(
     const size_t c = idx % op.chan_in;
     idx /= op.chan_in;
     const size_t o = idx % op.chan_out;
-    idx /= op.chan_out;
 
     auto i_tr = c * (op.chan_out * op.kernel * op.kernel) + o * (op.kernel * op.kernel) + k1 * (op.kernel) + k2;
     auto i_no = o * strides[0] + c * strides[1] + k1 * strides[2] + k2 * strides[3];
 
-    const T f = filters[i_no];
-    for (auto b = 0; b < op.batch; b++) {
-        filters_tr[b * numel + i_tr] = f;
-    }
+    filters_tr[i_tr] = filters[i_no];
 }
 
 template<typename T>
@@ -205,7 +181,7 @@ extern "C" __global__ void TR_FILTERS( \
     const size_t *strides, \
     TYPENAME *filters_tr \
 ) { \
-    transpose_and_broadcast_filters(op, filters, strides, filters_tr); \
+    transpose_filters(op, filters, strides, filters_tr); \
 } \
 extern "C" __global__ void SUM_TR_FILTERS( \
     const Conv2DOp op, \
@@ -220,13 +196,13 @@ CONV_OP(
     float,
     unfold_input_into_patches_f32,
     unfold_output_into_patches_f32,
-    transpose_and_broadcast_filters_f32,
+    transpose_filters_f32,
     sum_transposed_filters_f32
 );
 CONV_OP(
     double,
     unfold_input_into_patches_f64,
     unfold_output_into_patches_f64,
-    transpose_and_broadcast_filters_f64,
+    transpose_filters_f64,
     sum_transposed_filters_f64
 );
\ No newline at end of file
diff --git a/src/tensor_ops/conv2d/cpu_kernel.rs b/src/tensor_ops/conv2d/cpu_kernel.rs
index 32dac86d1..e9b6686d5 100644
--- a/src/tensor_ops/conv2d/cpu_kernel.rs
+++ b/src/tensor_ops/conv2d/cpu_kernel.rs
@@ -1,5 +1,5 @@
 use crate::shapes::{Dtype, Shape};
-use crate::tensor::{cpu::*, Tensor};
+use crate::tensor::{cpu::*, Tensor, ZerosTensor};
 use crate::tensor_ops::matmul::cpu_kernel::MatMulImpl;
 
 use super::{Conv2DKernel, Conv2DOp};
@@ -163,6 +163,10 @@ impl<E: Dtype> Conv2DKernel<E> for Cpu
 where
     Self: MatMulImpl<E>,
 {
+    fn alloc<S: Shape>(&self, s: S) -> Result<Tensor<S, E, Self>, Self::Err> {
+        self.try_zeros_like(&s)
+    }
+
     fn forward<L: Shape, R: Shape, O: Shape>(
         &self,
         op: Conv2DOp,
diff --git a/src/tensor_ops/conv2d/cuda_kernel.rs b/src/tensor_ops/conv2d/cuda_kernel.rs
index a1eb4b885..f6f7fa006 100644
--- a/src/tensor_ops/conv2d/cuda_kernel.rs
+++ b/src/tensor_ops/conv2d/cuda_kernel.rs
@@ -1,10 +1,10 @@
 use cudarc::cublas::{CudaBlas, Gemm};
-use cudarc::driver::{DeviceRepr, DeviceSlice, LaunchAsync, ValidAsZeroBits};
+use cudarc::driver::{DeviceRepr, LaunchAsync, ValidAsZeroBits};
 
 use crate::tensor_ops::matmul::cuda_kernel::sgemm_batch;
 use crate::{
     shapes::*,
-    tensor::{launch_cfg, Cuda, Tensor},
+    tensor::{launch_cfg, unique_id, Cuda, Tensor},
 };
 
 use std::sync::Arc;
@@ -23,7 +23,7 @@ impl HasCudaKernel<f32> for Cuda {
     const FNS: &'static [&'static str] = &[
         "unfold_input_into_patches_f32",
         "unfold_output_into_patches_f32",
-        "transpose_and_broadcast_filters_f32",
+        "transpose_filters_f32",
         "sum_transposed_filters_f32",
     ];
 }
@@ -33,7 +33,7 @@ impl HasCudaKernel<f64> for Cuda {
     const FNS: &'static [&'static str] = &[
         "unfold_input_into_patches_f64",
         "unfold_output_into_patches_f64",
-        "transpose_and_broadcast_filters_f64",
+        "transpose_filters_f64",
         "sum_transposed_filters_f64",
     ];
 }
@@ -51,6 +51,17 @@ where
     Self: HasCudaKernel<E>,
     CudaBlas: Gemm<E>,
 {
+    fn alloc<S: Shape>(&self, shape: S) -> Result<Tensor<S, E, Self>, Self::Err> {
+        let data = Arc::new(unsafe { self.dev.alloc::<E>(shape.num_elements()) }?);
+        Ok(Tensor {
+            id: unique_id(),
+            data,
+            shape,
+            strides: shape.strides(),
+            device: self.clone(),
+            tape: Default::default(),
+        })
+    }
     fn forward<L: Shape, R: Shape, O: Shape>(
         &self,
         op: super::Conv2DOp,
@@ -62,11 +73,15 @@ where
             self.dev.load_ptx(PTX_SRC.into(), Self::MOD, Self::FNS)?;
         }
 
-        let patches_numel = op.batch * op.chan_in * op.kernel * op.kernel * op.h_out * op.w_out;
-        let mut patches = self.dev.alloc_zeros::<E>(patches_numel)?;
+        let patches_item_numel = op.chan_in * op.kernel * op.kernel * op.h_out * op.w_out;
+        let patches_numel = op.batch * patches_item_numel;
+
+        let mut patches = unsafe { self.get_workspace::<E>(patches_numel) }?;
+        let mut patches = unsafe { patches.transmute_mut::<E>(patches_numel).unwrap() };
+
         let img_strides = self.dev.htod_copy(make_4d::<L>(lhs.strides).into())?;
         let unfold_fn = self.dev.get_func(Self::MOD, Self::FNS[0]).unwrap();
-        let cfg = launch_cfg(patches.len() as u32);
+        let cfg = launch_cfg(patches_numel as u32);
         let params = (op, lhs.data.as_ref(), &img_strides, &mut patches);
         unsafe { unfold_fn.launch(cfg, params) }?;
 
@@ -83,7 +98,7 @@ where
                 &patches,
                 [k * n, n, 1],
                 Default::default(),
-                Arc::make_mut(&mut out.data),
+                Arc::get_mut(&mut out.data).unwrap(),
                 [m * n, n, 1],
             )
             .unwrap();
@@ -102,41 +117,53 @@ where
         _: &Tensor<O, E, Self>,
         grad_out: &Self::Vec<E>,
     ) -> Result<(), Self::Err> {
-        let patches_numel = op.batch * op.chan_out * op.kernel * op.kernel * op.h_in * op.w_in;
-        let mut patches = self.dev.alloc_zeros::<E>(patches_numel)?;
+        let patches_item_numel = op.chan_out * op.kernel * op.kernel * op.h_in * op.w_in;
+        let patches_numel = op.batch * patches_item_numel;
+        let filters_numel = op.chan_in * op.chan_out * op.kernel * op.kernel;
+
+        let mut patches = unsafe { self.get_workspace::<E>(patches_numel) }?;
+        let mut patches = unsafe { patches.transmute_mut::<E>(patches_numel).unwrap() };
+
+        let mut f_b1023 = unsafe { self.dev.alloc::<E>(filters_numel) }?;
+        let mut grad_f_b1023 = unsafe { self.dev.alloc::<E>(op.batch * filters_numel) }?;
+        let f_strides = self.dev.htod_copy(rhs.strides.into())?;
+
+        self.par_stream.wait_for_default()?;
 
         {
             // unfold grad_out into patches
             let unfold_fn = self.dev.get_func(Self::MOD, Self::FNS[1]).unwrap();
-            let cfg = launch_cfg(patches_numel as u32);
+            let cfg = launch_cfg(patches_item_numel as u32);
             unsafe { unfold_fn.launch(cfg, (op, grad_out, &mut patches)) }?;
         }
 
-        let filters_numel = op.batch * op.chan_in * op.chan_out * op.kernel * op.kernel;
-        let mut f_b1023 = self.dev.alloc_zeros::<E>(filters_numel)?;
-        let mut grad_f_b1023 = self.dev.alloc_zeros::<E>(filters_numel)?;
-        let f_strides = self.dev.htod_copy(rhs.strides.into())?;
-
         {
             // prepare filters for backward operations by
-            // swapping dims 0 and 1 and adding a batch dimension
+            // swapping dims 0 and 1
             let tr_fn = self.dev.get_func(Self::MOD, Self::FNS[2]).unwrap();
             let cfg = launch_cfg(rhs.shape.num_elements() as u32);
-            unsafe { tr_fn.launch(cfg, (op, rhs.data.as_ref(), &f_strides, &mut f_b1023)) }?;
-        }
+            unsafe {
+                tr_fn.launch_on_stream(
+                    self.par_stream.as_ref(),
+                    cfg,
+                    (op, rhs.data.as_ref(), &f_strides, &mut f_b1023),
+                )
+            }?;
+
+            self.par_stream.wait_for_default()?;
 
-        {
             // img_g += filters * patches
             // (B, C, H * W) += (B, C, O * K * K) * (B, O * K * K, H * W)
             let m = op.chan_in;
             let k = op.chan_out * op.kernel * op.kernel;
             let n = op.h_in * op.w_in;
             unsafe {
+                self.blas.set_stream(Some(self.par_stream.as_ref()))?;
                 sgemm_batch(
                     self.blas.as_ref(),
                     (op.batch, m, k, n),
                     &f_b1023,
-                    [m * k, k, 1],
+                    [0, k, 1],
                     &patches,
                     [k * n, n, 1],
                     <E>::ONE,
@@ -144,6 +171,7 @@ where
                     [m * n, n, 1],
                 )
                 .unwrap();
+                self.blas.set_stream(None)?;
             }
         }
 
@@ -161,7 +189,7 @@ where
                     [m * k, k, 1],
                     &patches,
                     [k * n, 1, k],
-                    <E>::ONE,
+                    Default::default(),
                     &mut grad_f_b1023,
                     [m * n, n, 1],
                 )
@@ -175,6 +203,8 @@ where
             unsafe { sum_fn.launch(cfg, (op, &grad_f_b1023, grad_rhs, &f_strides)) }?;
         }
 
+        self.dev.wait_for(self.par_stream.as_ref())?;
+
         Ok(())
     }
 }
diff --git a/src/tensor_ops/conv2d/mod.rs b/src/tensor_ops/conv2d/mod.rs
index b496dec82..48ba7946e 100644
--- a/src/tensor_ops/conv2d/mod.rs
+++ b/src/tensor_ops/conv2d/mod.rs
@@ -55,6 +55,8 @@ impl Conv2DOp {
 }
 
 pub(super) trait Conv2DKernel<E: Dtype>: DeviceStorage {
+    fn alloc<S: Shape>(&self, s: S) -> Result<Tensor<S, E, Self>, Self::Err>;
+
     fn forward<L: Shape, R: Shape, O: Shape>(
         &self,
         op: Conv2DOp,
@@ -151,7 +153,9 @@ where
         let (lhs, ltape) = self.split_tape();
         let (rhs, rtape) = filters.split_tape();
         let mut tape = ltape.merge(rtape);
-        let mut out = lhs.device.try_zeros()?;
+        let mut out = lhs
+            .device
+            .alloc((Const, Default::default(), Default::default()))?;
         lhs.device.forward(op, &lhs, &rhs, &mut out)?;
         let phantom_out = out.clone();
         tape.try_alloc_grad(&lhs)?;
@@ -203,9 +207,9 @@ where
         let op = Conv2DOp::new(S, P, K, [batch.size(), C, H, W], O);
         let (lhs, ltape) = self.split_tape();
         let (rhs, rtape) = filters.split_tape();
-        let mut out =
-            lhs.device
-                .try_zeros_like(&(batch, Const, Default::default(), Default::default()))?;
+        let mut out = lhs
+            .device
+            .alloc((batch, Const, Default::default(), Default::default()))?;
         let mut tape = ltape.merge(rtape);
         lhs.device.forward(op, &lhs, &rhs, &mut out)?;
         let phantom_out = out.clone();
diff --git a/src/tensor_ops/utilities/cuda_kernels.rs b/src/tensor_ops/utilities/cuda_kernels.rs
index f1a180e64..16f702397 100644
--- a/src/tensor_ops/utilities/cuda_kernels.rs
+++ b/src/tensor_ops/utilities/cuda_kernels.rs
@@ -234,10 +234,22 @@ impl<E: Dtype, K: BinaryOpCudaKernel<E> + DeviceRepr + Clone> BinaryKernel<K, E>
             lhs.strides,
         );
 
+        let ((out_dims2, out_strides2), lhs_strides2) = permute_for_binary_backward(
+            lhs.shape.concrete(),
+            lhs.shape.strides(),
+            lhs.strides,
+            rhs.strides,
+        );
+
         let out_dims1 = self.dev.htod_copy(out_dims1)?;
         let out_strides1 = self.dev.htod_copy(out_strides1)?;
         let rhs_strides1 = self.dev.htod_copy(rhs_strides1)?;
+        let out_dims2 = self.dev.htod_copy(out_dims2)?;
+        let out_strides2 = self.dev.htod_copy(out_strides2)?;
+        let lhs_strides2 = self.dev.htod_copy(lhs_strides2)?;
+
         let chunk_len1 = numel / physical_numel(lhs.shape.concrete(), lhs.strides);
+        let chunk_len2 = numel / physical_numel(rhs.shape.concrete(), rhs.strides);
 
         let params_lhs = (
             op.clone(),        // const OP_STRUCT op,
@@ -256,18 +268,6 @@ impl<E: Dtype, K: BinaryOpCudaKernel<E> + DeviceRepr + Clone> BinaryKernel<K, E>
         self.par_stream.wait_for_default()?;
         unsafe { bwd_lhs_fn.launch_on_stream(&self.par_stream, cfg, params_lhs) }?;
 
-        let ((out_dims2, out_strides2), lhs_strides2) = permute_for_binary_backward(
-            lhs.shape.concrete(),
-            lhs.shape.strides(),
-            lhs.strides,
-            rhs.strides,
-        );
-
-        let out_dims2 = self.dev.htod_copy(out_dims2)?;
-        let out_strides2 = self.dev.htod_copy(out_strides2)?;
-        let lhs_strides2 = self.dev.htod_copy(lhs_strides2)?;
-        let chunk_len2 = numel / physical_numel(rhs.shape.concrete(), rhs.strides);
-
         let params_rhs = (
             op,                // const OP_STRUCT op,
             numel,             // const size_t numel,


Test patch: