Fixes sub-array (cpu, cuda, opencl, oneapi) support for convolve

willyborn · Jun 10, 2025 · Jun 27, 2025 · Jun 27, 2025 · Jun 27, 2025 · Jun 27, 2025
commit a3a91e1958591f8d0d782d13248a67df6a3d3ae5
diff --git a/src/backend/common/KernelInterface.hpp b/src/backend/common/KernelInterface.hpp
@@ -64,6 +64,51 @@ class KernelInterface {
     virtual void copyToReadOnly(DevPtrType dst, DevPtrType src,
                                 size_t bytes) = 0;
 
+    /// \brief Copy data from device memory to read-only memory
+    ///
+    /// This function copies data of `bytes` size from the device pointer to a
+    /// read-only memory.
+    ///
+    /// \param[in] dst is the device pointer to which data will be copied
+    /// \param[in] src is the device pointer from which data will be copied
+    /// \param[in] srcXInBytes is offset in Bytes
+    /// \param[in] bytes are the number of bytes of data to be copied
+    virtual void copyToReadOnly(DevPtrType dst, DevPtrType src,
+                                size_t srcXInBytes, size_t bytes) = 0;
+
+    /// \brief Copy strided 2D data from device memory to read-only memory
+    ///
+    /// This function copies data of any 2D array from the device pointer to a
+    /// read-only memory.
+    ///
+    /// \param[in] dst is the device pointer to which data will be copied
+    /// \param[in] src is the device pointer from which data will be copied
+    /// \param[in] srcXInBytes is offset in Bytes
+    /// \param[in] srcPitchInBytes is strides[1] in Bytes
+    /// \param[in] height is the number of elements for dim[1] dst
+    /// \param[in] widthInBytes are #bytes of continous data to copy (dim[0])
+    virtual void copyToReadOnly2D(DevPtrType dst, DevPtrType src,
+                                  size_t srcXInBytes, size_t srcPitchInBytes,
+                                  size_t height, size_t widthInBytes) = 0;
+
+    /// \brief Copy strided 3D data from device memory to read-only memory
+    ///
+    /// This function copies data of any 3D array from the device pointer to a
+    /// read-only memory.
+    ///
+    /// \param[in] dst is the device pointer to which data will be copied
+    /// \param[in] src is the device pointer from which data will be copied
+    /// \param[in] srcXInBytes is offset in Bytes
+    /// \param[in] srcPitchInBytes is strides[1] in Bytes
+    /// \param[in] srcHeight is the number of elements ALLOCATED for dim[1] src
+    /// \param[in] depth is the number of elements for dim[2] dst
+    /// \param[in] height is the number of elements for dim[1] dst
+    /// \param[in] widthInBytes are #bytes of continous data to copy (dim[0])
+    virtual void copyToReadOnly3D(DevPtrType dst, DevPtrType src,
+                                  size_t srcXInBytes, size_t srcPitchInBytes,
+                                  size_t srcHeight, size_t depth, size_t height,
+                                  size_t widthInBytes) = 0;
+
     /// \brief Copy a single scalar to device memory
     ///
     /// This function copies a single value of type T from host variable

diff --git a/src/backend/cpu/kernel/convolve.hpp b/src/backend/cpu/kernel/convolve.hpp
@@ -125,8 +125,8 @@ void one2one_3d(InT *optr, InT const *const iptr, AccT const *const fptr,
                 }
                 optr[koff + joff + i - iStart] = InT(accum);
             }  // i loop ends here
-        }      // j loop ends here
-    }          // k loop ends here
+        }  // j loop ends here
+    }  // k loop ends here
 }
 
 template<typename InT, typename AccT>
@@ -217,7 +217,6 @@ void convolve2_separable(InT *optr, InT const *const iptr,
                          dim_t fDim, af::dim4 const &oStrides,
                          af::dim4 const &sStrides, dim_t fStride) {
     UNUSED(orgDims);
-    UNUSED(sStrides);
     UNUSED(fStride);
     for (dim_t j = 0; j < oDims[1]; ++j) {
         dim_t jOff = j * oStrides[1];
@@ -237,14 +236,18 @@ void convolve2_separable(InT *optr, InT const *const iptr,
                     dim_t offi     = ci - f;
                     bool isCIValid = offi >= 0 && offi < sDims[0];
                     bool isCJValid = cj >= 0 && cj < sDims[1];
-                    s_val = (isCJValid && isCIValid ? iptr[cj * sDims[0] + offi]
-                                                    : scalar<InT>(0));
+                    s_val =
+                        (isCJValid && isCIValid ? iptr[cj * sStrides.dims[1] +
+                                                       offi * sStrides.dims[0]]
+                                                : scalar<InT>(0));
                 } else {
                     dim_t offj     = cj - f;
                     bool isCIValid = ci >= 0 && ci < sDims[0];
                     bool isCJValid = offj >= 0 && offj < sDims[1];
-                    s_val = (isCJValid && isCIValid ? iptr[offj * sDims[0] + ci]
-                                                    : scalar<InT>(0));
+                    s_val =
+                        (isCJValid && isCIValid ? iptr[offj * sStrides.dims[1] +
+                                                       ci * sStrides.dims[0]]
+                                                : scalar<InT>(0));
                 }
 
                 accum += AccT(s_val * f_val);

diff --git a/src/backend/cuda/Kernel.cpp b/src/backend/cuda/Kernel.cpp
@@ -26,6 +26,67 @@ void Kernel::copyToReadOnly(Kernel::DevPtrType dst, Kernel::DevPtrType src,
     CU_CHECK(cuMemcpyDtoDAsync(dst, src, bytes, getActiveStream()));
 }
 
+void Kernel::copyToReadOnly(Kernel::DevPtrType dst, Kernel::DevPtrType src,
+                            size_t srcXInBytes, size_t bytes) {
+    CU_CHECK(cuMemcpyDtoDAsync(dst, src, bytes, getActiveStream()));
+}
+
+void Kernel::copyToReadOnly2D(Kernel::DevPtrType dst, Kernel::DevPtrType src,
+                              size_t srcXInBytes, size_t srcPitchInBytes,
+                              size_t height, size_t widthInBytes) {
+    CUDA_MEMCPY2D pCopy;
+    pCopy.srcXInBytes   = srcXInBytes;
+    pCopy.srcY          = 0;
+    pCopy.srcMemoryType = CU_MEMORYTYPE_DEVICE;
+    pCopy.srcDevice     = src;
+    pCopy.srcPitch      = srcPitchInBytes;
+
+    pCopy.dstXInBytes   = 0;
+    pCopy.dstY          = 0;
+    pCopy.dstMemoryType = CU_MEMORYTYPE_DEVICE;
+    pCopy.dstDevice     = dst;
+    pCopy.dstPitch      = widthInBytes;
+
+    pCopy.WidthInBytes = widthInBytes;
+    pCopy.Height       = height;
+    // CUdeviceptr srcStart = srcDevice + srcY*srcPitch + srcXInBytes;
+    // CUdeviceptr dstStart = dstDevice + dstY*dstPitch + dstXInBytes;
+
+    CU_CHECK(cuMemcpy2DAsync(&pCopy, getActiveStream()));
+}
+
+void Kernel::copyToReadOnly3D(Kernel::DevPtrType dst, Kernel::DevPtrType src,
+                              size_t srcXInBytes, size_t srcPitchInBytes,
+                              size_t srcHeight, size_t depth, size_t height,
+                              size_t widthInBytes) {
+    CUDA_MEMCPY3D pCopy;
+    pCopy.srcXInBytes   = srcXInBytes;
+    pCopy.srcY          = 0;
+    pCopy.srcZ          = 0;
+    pCopy.srcLOD        = 0;
+    pCopy.srcMemoryType = CU_MEMORYTYPE_DEVICE;
+    pCopy.srcDevice     = src;
+    pCopy.srcPitch      = srcPitchInBytes;
+    pCopy.srcHeight     = srcHeight;
+
+    pCopy.dstXInBytes   = 0;
+    pCopy.dstY          = 0;
+    pCopy.dstZ          = 0;
+    pCopy.dstMemoryType = CU_MEMORYTYPE_DEVICE;
+    pCopy.dstDevice     = dst;
+    pCopy.dstPitch      = widthInBytes;
+    pCopy.dstHeight     = height;
+
+    pCopy.WidthInBytes = widthInBytes;
+    pCopy.Height       = height;
+    pCopy.Depth        = depth;
+    // CUdeviceptr srcStart =
+    //      srcDevice + (srcZ*srcHeight+srcY)*srcPitch + srcXInBytes;
+    // CUdeviceptr dstStart =
+    //      dstDevice + (dstZ*dstHeight+dstY)*dstPitch + dstXInBytes;
+    CU_CHECK(cuMemcpy3DAsync(&pCopy, getActiveStream()));
+}
+
 void Kernel::setFlag(Kernel::DevPtrType dst, int* scalarValPtr,
                      const bool syncCopy) {
     CU_CHECK(

diff --git a/src/backend/cuda/Kernel.hpp b/src/backend/cuda/Kernel.hpp
@@ -66,6 +66,18 @@ class Kernel
 
     void copyToReadOnly(DevPtrType dst, DevPtrType src, size_t bytes) final;
 
+    void copyToReadOnly(DevPtrType dst, DevPtrType src, size_t srcXInBytes,
+                        size_t bytes) final;
+
+    void copyToReadOnly2D(DevPtrType dst, DevPtrType src, size_t srcXInBytes,
+                          size_t srcPitchInBytes, size_t height,
+                          size_t widthInBytes) final;
+
+    void copyToReadOnly3D(DevPtrType dst, DevPtrType src, size_t srcXInBytes,
+                          size_t srcPitchInBytes, size_t srcHeight,
+                          size_t depth, size_t height,
+                          size_t widthInBytes) final;
+
     void setFlag(DevPtrType dst, int* scalarValPtr,
                  const bool syncCopy = false) final;
 

diff --git a/src/backend/cuda/kernel/convolve.hpp b/src/backend/cuda/kernel/convolve.hpp
@@ -120,7 +120,6 @@ void convolve_1d(conv_kparam_t& p, Param<T> out, CParam<T> sig, CParam<aT> filt,
                 int f1Off      = b1 * filt.strides[1];
                 const aT* fptr = filt.ptr + (f1Off + f2Off + f3Off);
 
-                // FIXME: case where filter array is strided
                 auto constMemPtr = convolve1.getDevPtr(conv_c_name);
                 convolve1.copyToReadOnly(constMemPtr,
                                          reinterpret_cast<CUdeviceptr>(fptr),
@@ -145,28 +144,40 @@ void convolve_1d(conv_kparam_t& p, Param<T> out, CParam<T> sig, CParam<aT> filt,
 
 template<typename T, typename aT>
 void conv2Helper(const conv_kparam_t& p, Param<T> out, CParam<T> sig,
-                 const aT* fptr, int f0, int f1, const bool expand) {
-    const bool isFilterSizeLt5  = (f0 <= 5 && f1 <= 5);
-    const bool isFilterGt5AndSq = (f0 == f1 && f0 > 5 && f0 < 18);
+                 const aT* fptr, const dim_t* const fdims,
+                 const dim_t* const fstrides, const bool expand) {
+    const bool isFilterSizeLt5 = (fdims[0] <= 5 && fdims[1] <= 5);
+    const bool isFilterGt5AndSq =
+        (fdims[0] == fdims[1] && fdims[0] > 5 && fdims[0] < 18);
 
     if (!(isFilterSizeLt5 || isFilterGt5AndSq)) {
         char errMessage[256];
         snprintf(errMessage, sizeof(errMessage),
-                 "\nCUDA Convolution doesn't support %dx%d kernel\n", f0, f1);
+                 "\nCUDA Convolution doesn't support %lldx%lld kernel\n",
+                 fdims[0], fdims[1]);
         CUDA_NOT_SUPPORTED(errMessage);
     }
 
     auto convolve2 = common::getKernel(
         "arrayfire::cuda::convolve2", {{convolve2_cuh_src}},
         TemplateArgs(TemplateTypename<T>(), TemplateTypename<aT>(),
-                     TemplateArg(expand), TemplateArg(f0), TemplateArg(f1)),
+                     TemplateArg(expand), TemplateArg(fdims[0]),
+                     TemplateArg(fdims[1])),
         {{DefineValue(MAX_CONV1_FILTER_LEN), DefineValue(CONV_THREADS),
           DefineValue(CONV2_THREADS_X), DefineValue(CONV2_THREADS_Y)}});
 
-    // FIXME: case where filter array is strided
     auto constMemPtr = convolve2.getDevPtr(conv_c_name);
-    convolve2.copyToReadOnly(constMemPtr, reinterpret_cast<CUdeviceptr>(fptr),
-                             f0 * f1 * sizeof(aT));
+    if (fstrides[1] == fdims[0]) {
+        // linear filter array
+        convolve2.copyToReadOnly(constMemPtr,
+                                 reinterpret_cast<CUdeviceptr>(fptr),
+                                 fdims[0] * fdims[1] * sizeof(aT));
+    } else {
+        // strided filter array
+        convolve2.copyToReadOnly2D(
+            constMemPtr, reinterpret_cast<CUdeviceptr>(fptr), 0,
+            fstrides[1] * sizeof(aT), fdims[1], fdims[0] * sizeof(aT));
+    }
 
     EnqueueArgs qArgs(p.mBlocks, p.mThreads, getActiveStream());
     convolve2(qArgs, out, sig, p.mBlk_x, p.mBlk_y, p.o[1], p.o[2], p.s[1],
@@ -192,7 +203,7 @@ void convolve_2d(conv_kparam_t& p, Param<T> out, CParam<T> sig, CParam<aT> filt,
             p.s[1] = (p.inHasNoOffset ? 0 : b2);
             p.s[2] = (p.inHasNoOffset ? 0 : b3);
 
-            conv2Helper<T, aT>(p, out, sig, fptr, filt.dims[0], filt.dims[1],
+            conv2Helper<T, aT>(p, out, sig, fptr, filt.dims, filt.strides,
                                expand);
         }
     }
@@ -218,11 +229,18 @@ void convolve_3d(conv_kparam_t& p, Param<T> out, CParam<T> sig, CParam<aT> filt,
 
         const aT* fptr = filt.ptr + f3Off;
 
-        // FIXME: case where filter array is strided
         auto constMemPtr = convolve3.getDevPtr(conv_c_name);
-        convolve3.copyToReadOnly(
-            constMemPtr, reinterpret_cast<CUdeviceptr>(fptr), filterSize);
-
+        if (filt.strides[2] == filt.dims[0] * filt.dims[1]) {
+            // linear filter array
+            convolve3.copyToReadOnly(
+                constMemPtr, reinterpret_cast<CUdeviceptr>(fptr), filterSize);
+        } else {
+            // strided filter array
+            convolve3.copyToReadOnly3D(
+                constMemPtr, reinterpret_cast<CUdeviceptr>(fptr), 0,
+                filt.strides[1] * sizeof(aT), filt.strides[2] / filt.strides[1],
+                filt.dims[2], filt.dims[1], filt.dims[0] * sizeof(aT));
+        }
         p.o[2] = (p.outHasNoOffset ? 0 : b3);
         p.s[2] = (p.inHasNoOffset ? 0 : b3);
 
@@ -321,11 +339,26 @@ void convolve2(Param<T> out, CParam<T> signal, CParam<aT> filter, int conv_dim,
 
     dim3 blocks(blk_x * signal.dims[2], blk_y * signal.dims[3]);
 
-    // FIXME: case where filter array is strided
     auto constMemPtr = convolve2_separable.getDevPtr(sconv_c_name);
-    convolve2_separable.copyToReadOnly(
-        constMemPtr, reinterpret_cast<CUdeviceptr>(filter.ptr),
-        fLen * sizeof(aT));
+    if (filter.strides[1] == filter.dims[0]) {
+        // linear filter array
+        convolve2_separable.copyToReadOnly(
+            constMemPtr, reinterpret_cast<CUdeviceptr>(filter.ptr),
+            fLen * sizeof(aT));
+    } else {
+        // strided filter array 4D
+        const aT* fptr   = filter.ptr;
+        CUdeviceptr dptr = constMemPtr;
+        for (dim_t d3 = 0; d3 < filter.dims[3]; ++d3, fptr += filter.strides[3],
+                   dptr += filter.dims[0] * filter.dims[1] * filter.dims[2] *
+                                                sizeof(aT)) {
+            convolve2_separable.copyToReadOnly3D(
+                dptr, reinterpret_cast<CUdeviceptr>(fptr), 0,
+                filter.strides[1] * sizeof(aT),
+                filter.strides[2] + filter.strides[1], filter.dims[2],
+                filter.dims[1], filter.dims[0] * sizeof(aT));
+        }
+    }
 
     EnqueueArgs qArgs(blocks, threads, getActiveStream());
     convolve2_separable(qArgs, out, signal, blk_x, blk_y);

diff --git a/src/backend/oneapi/kernel/convolve.hpp b/src/backend/oneapi/kernel/convolve.hpp
@@ -13,6 +13,7 @@
 #include <common/kernel_cache.hpp>
 #include <debug_oneapi.hpp>
 #include <kernel/accessors.hpp>
+#include <kernel/memcopy.hpp>
 #include <af/defines.h>
 
 #include <sycl/sycl.hpp>
@@ -75,7 +76,7 @@ void prepareKernelArgs(conv_kparam_t<aT> &param, dim_t *oDims,
         param.nBBS0    = divup(oDims[0], THREADS);
         param.nBBS1    = batchDims[2];
         param.global   = range<3>(param.nBBS0 * THREADS * batchDims[1],
-                                param.nBBS1 * batchDims[3], 1);
+                                  param.nBBS1 * batchDims[3], 1);
         param.loc_size = (THREADS + 2 * (fDims[0] - 1));
     } else if (rank == 2) {
         param.local  = range<3>{THREADS_X, THREADS_Y, 1};
@@ -89,26 +90,13 @@ void prepareKernelArgs(conv_kparam_t<aT> &param, dim_t *oDims,
         param.nBBS1    = divup(oDims[1], CUBE_Y);
         int blk_z      = divup(oDims[2], CUBE_Z);
         param.global   = range<3>(param.nBBS0 * CUBE_X * batchDims[3],
-                                param.nBBS1 * CUBE_Y, blk_z * CUBE_Z);
+                                  param.nBBS1 * CUBE_Y, blk_z * CUBE_Z);
         param.loc_size = (CUBE_X + 2 * (fDims[0] - 1)) *
                          (CUBE_Y + 2 * (fDims[1] - 1)) *
                          (CUBE_Z + 2 * (fDims[2] - 1));
     }
 }
 
-template<typename T>
-void memcpyBuffer(sycl::buffer<T, 1> &dest, sycl::buffer<T, 1> &src,
-                  const size_t n, const size_t srcOffset) {
-    getQueue().submit([&](auto &h) {
-        sycl::accessor srcAcc{src, h, sycl::range{n}, sycl::id{srcOffset},
-                              sycl::read_only};
-        sycl::accessor destAcc{
-            dest,         h, sycl::range{n}, sycl::id{0}, sycl::write_only,
-            sycl::no_init};
-        h.copy(srcAcc, destAcc);
-    });
-}
-
 #include "convolve1.hpp"
 #include "convolve2.hpp"
 #include "convolve3.hpp"

diff --git a/src/backend/oneapi/kernel/convolve1.hpp b/src/backend/oneapi/kernel/convolve1.hpp
@@ -134,20 +134,23 @@ void conv1Helper(const conv_kparam_t<aT> &param, Param<T> &out,
 template<typename T, typename aT>
 void conv1(conv_kparam_t<aT> &p, Param<T> &out, const Param<T> &sig,
            const Param<aT> &filt, const bool expand) {
-    const size_t se_size = filt.info.dims[0];
-    sycl::buffer<aT> impulse{sycl::range(filt.info.dims[0])};
-    int f0Off = filt.info.offset;
+    const dim_t se_size = filt.info.dims[0];
+    sycl::buffer<aT> impulse{sycl::range(se_size)};
+    const dim_t mstrides[4] = {1, se_size, se_size, se_size};
+    const dim_t mdims[4]    = {filt.info.dims[0], 1, 1, 1};
+    const dim_t f0Off       = filt.info.offset;
     for (int b3 = 0; b3 < filt.info.dims[3]; ++b3) {
-        int f3Off = b3 * filt.info.strides[3];
+        const dim_t f3Off = b3 * filt.info.strides[3];
 
         for (int b2 = 0; b2 < filt.info.dims[2]; ++b2) {
-            int f2Off = b2 * filt.info.strides[2];
+            const dim_t f2Off = b2 * filt.info.strides[2];
 
             for (int b1 = 0; b1 < filt.info.dims[1]; ++b1) {
-                int f1Off = b1 * filt.info.strides[1];
+                const dim_t f1Off = b1 * filt.info.strides[1];
 
-                const size_t srcOffset = f0Off + f1Off + f2Off + f3Off;
-                memcpyBuffer(impulse, *filt.data, se_size, srcOffset);
+                const dim_t srcOffset = f0Off + f1Off + f2Off + f3Off;
+                kernel::memcopy(&impulse, mstrides, filt.data, mdims,
+                                filt.info.strides, srcOffset, 1);
                 p.impulse = &impulse;
 
                 p.o[0] = (p.outHasNoOffset ? 0 : b1);