diff --git a/modules/core/include/opencv2/core/cuda.hpp b/modules/core/include/opencv2/core/cuda.hpp
index 15f0416971..31768e0686 100644
--- a/modules/core/include/opencv2/core/cuda.hpp
+++ b/modules/core/include/opencv2/core/cuda.hpp
@@ -1335,6 +1335,16 @@ private:
 CV_EXPORTS_W void printCudaDeviceInfo(int device);
 CV_EXPORTS_W void printShortCudaDeviceInfo(int device);
 
+/** @brief Returns the MatAllocator that backs cv::UMat with CUDA device memory.
+
+Bind it to a UMat before allocation to place the buffer on the current CUDA device:
+@code
+    UMat u; u.allocator = cv::cuda::getCudaAllocator(); u.create(rows, cols, CV_32F);
+@endcode
+Returns NULL when OpenCV is built without CUDA. Mirrors cv::ocl::getOpenCLAllocator().
+*/
+CV_EXPORTS MatAllocator* getCudaAllocator();
+
 //! @} cudacore_init
 
 }} // namespace cv { namespace cuda {
diff --git a/modules/core/src/cuda_umat.cpp b/modules/core/src/cuda_umat.cpp
new file mode 100644
index 0000000000..ff5ccf2a55
--- /dev/null
+++ b/modules/core/src/cuda_umat.cpp
@@ -0,0 +1,188 @@
+// This file is part of OpenCV project.
+// It is subject to the license terms in the LICENSE file found in the top-level directory
+// of this distribution and at http://opencv.org/license.html.
+// Copyright (C) 2026, BigVision LLC, all rights reserved.
+// Third party copyrights are property of their respective owners.
+
+// CUDA-backed MatAllocator for UMat (Phase P1): synchronous host<->device transfers.
+
+#include "precomp.hpp"
+
+namespace cv { namespace cuda {
+
+#ifndef HAVE_CUDA
+
+MatAllocator* getCudaAllocator()
+{
+    return NULL;
+}
+
+#else
+
+class CudaUMatAllocator CV_FINAL : public MatAllocator
+{
+public:
+    UMatData* allocate(int dims, const int* sizes, int type, void* data0,
+                       size_t* step, AccessFlag /*flags*/, UMatUsageFlags /*usageFlags*/) const CV_OVERRIDE
+    {
+        const size_t elemSize = CV_ELEM_SIZE(type);
+        size_t total = elemSize;
+        for (int i = 0; i < dims; i++)
+            total *= (size_t)sizes[i];
+
+        if (step)
+        {
+            step[dims - 1] = elemSize;
+            for (int i = dims - 2; i >= 0; i--)
+                step[i] = step[i + 1] * (size_t)sizes[i + 1];
+        }
+
+        UMatData* u = new UMatData(this);
+        u->size = total;
+
+        if (total > 0)
+            cudaSafeCall(cudaMalloc(&u->handle, total));
+
+        if (data0)
+        {
+            u->data = u->origdata = static_cast<uchar*>(data0);
+            u->flags |= UMatData::USER_ALLOCATED;
+            u->markDeviceCopyObsolete(true);
+        }
+        else
+        {
+            u->markHostCopyObsolete(true);
+        }
+        return u;
+    }
+
+    bool allocate(UMatData* u, AccessFlag /*accessFlags*/, UMatUsageFlags /*usageFlags*/) const CV_OVERRIDE
+    {
+        if (!u)
+            return false;
+        if (!u->handle && u->size > 0)
+        {
+            cudaSafeCall(cudaMalloc(&u->handle, u->size));
+            if (u->data)
+                u->markDeviceCopyObsolete(true);
+            else
+                u->markHostCopyObsolete(true);
+        }
+        return true;
+    }
+
+    void deallocate(UMatData* u) const CV_OVERRIDE
+    {
+        if (!u)
+            return;
+        if (u->handle)
+            cudaSafeCall(cudaFree(u->handle));
+        if (u->data && !(u->flags & UMatData::USER_ALLOCATED))
+            fastFree(u->data);
+        delete u;
+    }
+
+    void map(UMatData* u, AccessFlag accessFlags) const CV_OVERRIDE
+    {
+        if (!u)
+            return;
+        if (!u->data && u->size > 0)
+            u->data = u->origdata = static_cast<uchar*>(fastMalloc(u->size));
+        if (u->hostCopyObsolete() && u->handle && u->data)
+        {
+            cudaSafeCall(cudaMemcpy(u->data, u->handle, u->size, cudaMemcpyDeviceToHost));
+            u->markHostCopyObsolete(false);
+        }
+        if (!!(accessFlags & ACCESS_WRITE))
+            u->markDeviceCopyObsolete(true);
+    }
+
+    void unmap(UMatData* u) const CV_OVERRIDE
+    {
+        if (!u)
+            return;
+        if (u->deviceCopyObsolete() && u->handle && u->data)
+        {
+            cudaSafeCall(cudaMemcpy(u->handle, u->data, u->size, cudaMemcpyHostToDevice));
+            u->markDeviceCopyObsolete(false);
+        }
+    }
+
+    void download(UMatData* u, void* dstptr, int dims, const size_t sz[],
+                  const size_t srcofs[], const size_t srcstep[],
+                  const size_t dststep[]) const CV_OVERRIDE
+    {
+        if (!u || !u->handle || !dstptr)
+            return;
+        copyPlanes((uchar*)u->handle, srcofs, srcstep, (uchar*)dstptr, 0, dststep,
+                   dims, sz, cudaMemcpyDeviceToHost);
+    }
+
+    void upload(UMatData* u, const void* srcptr, int dims, const size_t sz[],
+                const size_t dstofs[], const size_t dststep[],
+                const size_t srcstep[]) const CV_OVERRIDE
+    {
+        if (!u || !srcptr)
+            return;
+        if (!u->handle && u->size > 0)
+            cudaSafeCall(cudaMalloc(&u->handle, u->size));
+        copyPlanes((uchar*)srcptr, NULL, srcstep, (uchar*)u->handle, dstofs, dststep,
+                   dims, sz, cudaMemcpyHostToDevice);
+        u->markHostCopyObsolete(true);
+        u->markDeviceCopyObsolete(false);
+    }
+
+    // Same-world D2D only; UMat::copyTo calls copy() solely when both share this allocator.
+    void copy(UMatData* src, UMatData* dst, int dims, const size_t sz[],
+              const size_t srcofs[], const size_t srcstep[],
+              const size_t dstofs[], const size_t dststep[], bool /*sync*/) const CV_OVERRIDE
+    {
+        if (!src || !dst || !src->handle || !dst->handle)
+            return;
+        copyPlanes((uchar*)src->handle, srcofs, srcstep, (uchar*)dst->handle, dstofs, dststep,
+                   dims, sz, cudaMemcpyDeviceToDevice);
+        dst->markHostCopyObsolete(true);
+        dst->markDeviceCopyObsolete(false);
+    }
+
+private:
+    // sz/steps follow the MatAllocator convention: the last dim carries element bytes.
+    static void copyPlanes(uchar* srcbase, const size_t srcofs[], const size_t srcstep[],
+                           uchar* dstbase, const size_t dstofs[], const size_t dststep[],
+                           int dims, const size_t sz[], cudaMemcpyKind kind)
+    {
+        int isz[CV_MAX_DIM];
+        uchar* srcptr = srcbase;
+        uchar* dstptr = dstbase;
+        for (int i = 0; i < dims; i++)
+        {
+            CV_Assert(sz[i] <= (size_t)INT_MAX);
+            if (sz[i] == 0)
+                return;
+            if (srcofs)
+                srcptr += srcofs[i] * (i <= dims - 2 ? srcstep[i] : 1);
+            if (dstofs)
+                dstptr += dstofs[i] * (i <= dims - 2 ? dststep[i] : 1);
+            isz[i] = (int)sz[i];
+        }
+
+        Mat src(dims, isz, CV_8U, srcptr, srcstep);
+        Mat dst(dims, isz, CV_8U, dstptr, dststep);
+
+        const Mat* arrays[] = { &src, &dst };
+        uchar* ptrs[2];
+        NAryMatIterator it(arrays, ptrs, 2);
+        const size_t planesz = it.size;
+        for (size_t j = 0; j < it.nplanes; j++, ++it)
+            cudaSafeCall(cudaMemcpy(ptrs[1], ptrs[0], planesz, kind));
+    }
+};
+
+MatAllocator* getCudaAllocator()
+{
+    CV_SINGLETON_LAZY_INIT(CudaUMatAllocator, new CudaUMatAllocator())
+}
+
+#endif
+
+}} // namespace cv { namespace cuda {
diff --git a/modules/core/test/test_cuda_umat.cpp b/modules/core/test/test_cuda_umat.cpp
new file mode 100644
index 0000000000..8445cde9d2
--- /dev/null
+++ b/modules/core/test/test_cuda_umat.cpp
@@ -0,0 +1,231 @@
+// This file is part of OpenCV project.
+// It is subject to the license terms in the LICENSE file found in the top-level directory
+// of this distribution and at http://opencv.org/license.html.
+// Copyright (C) 2026, BigVision LLC, all rights reserved.
+// Third party copyrights are property of their respective owners.
+
+#include "test_precomp.hpp"
+
+#ifdef HAVE_CUDA
+
+#include <opencv2/core/cuda.hpp>
+#include <opencv2/core/ocl.hpp>
+
+namespace opencv_test { namespace {
+
+static void skipIfNoCudaDevice()
+{
+    if (cv::cuda::getCudaEnabledDeviceCount() <= 0)
+        throw SkipTestException("No CUDA device available");
+}
+
+static UMat makeCudaUMat()
+{
+    UMat u;
+    u.allocator = cv::cuda::getCudaAllocator();
+    return u;
+}
+
+TEST(Core_CudaUMat, roundtrip_contiguous)
+{
+    skipIfNoCudaDevice();
+    ASSERT_TRUE(cv::cuda::getCudaAllocator() != NULL);
+
+    Mat src(64, 48, CV_32FC1);
+    randu(src, 0, 1);
+
+    UMat u = makeCudaUMat();
+    src.copyTo(u);
+    ASSERT_TRUE(u.u != NULL && u.u->handle != NULL);
+
+    Mat dst;
+    u.copyTo(dst);
+
+    EXPECT_EQ(src.size(), dst.size());
+    EXPECT_EQ(src.type(), dst.type());
+    EXPECT_EQ(0.0, cvtest::norm(src, dst, NORM_INF));
+}
+
+TEST(Core_CudaUMat, roundtrip_multichannel)
+{
+    skipIfNoCudaDevice();
+
+    Mat src(37, 53, CV_8UC3);
+    randu(src, 0, 255);
+
+    UMat u = makeCudaUMat();
+    src.copyTo(u);
+
+    Mat dst;
+    u.copyTo(dst);
+    EXPECT_EQ(0.0, cvtest::norm(src, dst, NORM_INF));
+}
+
+TEST(Core_CudaUMat, roundtrip_roi_source)
+{
+    skipIfNoCudaDevice();
+
+    Mat full(100, 100, CV_32FC1);
+    randu(full, 0, 1);
+    Mat roi = full(Rect(10, 12, 50, 40));
+    ASSERT_FALSE(roi.isContinuous());
+
+    UMat u = makeCudaUMat();
+    roi.copyTo(u);
+
+    Mat dst;
+    u.copyTo(dst);
+    EXPECT_EQ(roi.size(), dst.size());
+    EXPECT_EQ(0.0, cvtest::norm(roi, dst, NORM_INF));
+}
+
+TEST(Core_CudaUMat, download_into_roi)
+{
+    skipIfNoCudaDevice();
+
+    Mat src(40, 50, CV_32FC1);
+    randu(src, 0, 1);
+
+    UMat u = makeCudaUMat();
+    src.copyTo(u);
+
+    Mat canvas(80, 90, CV_32FC1, Scalar(0));
+    Mat dstRoi = canvas(Rect(5, 7, 50, 40));
+    ASSERT_FALSE(dstRoi.isContinuous());
+    u.copyTo(dstRoi);
+
+    EXPECT_EQ(0.0, cvtest::norm(src, dstRoi, NORM_INF));
+}
+
+TEST(Core_CudaUMat, device_to_device_copy)
+{
+    skipIfNoCudaDevice();
+
+    Mat src(70, 90, CV_32FC1);
+    randu(src, 0, 1);
+
+    UMat a = makeCudaUMat();
+    src.copyTo(a);
+
+    UMat b = makeCudaUMat();
+    a.copyTo(b);
+
+    Mat dst;
+    b.copyTo(dst);
+    EXPECT_EQ(0.0, cvtest::norm(src, dst, NORM_INF));
+}
+
+// Proves the allocator's handle is a real, kernel-usable CUDA buffer via GpuMat interop.
+TEST(Core_CudaUMat, kernel_interop_via_gpumat)
+{
+    skipIfNoCudaDevice();
+
+    Mat src(32, 40, CV_32FC1, Scalar(1.0f));
+
+    UMat u = makeCudaUMat();
+    src.copyTo(u);
+    ASSERT_TRUE(u.u != NULL && u.u->handle != NULL);
+
+    cv::cuda::GpuMat g(u.rows, u.cols, u.type(), u.u->handle,
+                       (size_t)u.cols * CV_ELEM_SIZE(u.type()));
+
+    Mat viaGpu;
+    g.download(viaGpu);
+    EXPECT_EQ(0.0, cvtest::norm(src, viaGpu, NORM_INF));
+
+    g.setTo(Scalar(7.0f));
+    u.u->markHostCopyObsolete(true);          // device was mutated behind UMat's back
+    Mat afterKernel;
+    u.copyTo(afterKernel);
+    EXPECT_EQ(0.0, cvtest::norm(Mat(src.size(), src.type(), Scalar(7.0f)), afterKernel, NORM_INF));
+}
+
+// P1 gives a UMat storage on CUDA, not GPU execution. With OpenCL on, cv:: ops feed the
+// CUDA pointer to the driver as a cl_mem and segfault; the tests below force OpenCL off so
+// ops resolve through the host getMat() fallback, and pin why the OpenCL path is unusable.
+
+struct NoOpenCL
+{
+    bool prev;
+    NoOpenCL() : prev(cv::ocl::useOpenCL()) { cv::ocl::setUseOpenCL(false); }
+    ~NoOpenCL() { cv::ocl::setUseOpenCL(prev); }
+};
+
+// A CUDA UMat fails the OpenCL kernel path's isCLBuffer precondition (carries the CUDA allocator).
+TEST(Core_CudaUMat, is_not_opencl_buffer)
+{
+    skipIfNoCudaDevice();
+    UMat u = makeCudaUMat();
+    Mat(8, 8, CV_32FC1, Scalar(1)).copyTo(u);
+    ASSERT_TRUE(u.u != NULL);
+    EXPECT_EQ(cv::cuda::getCudaAllocator(), u.u->currAllocator);
+#ifdef HAVE_OPENCL
+    EXPECT_NE((MatAllocator*)cv::ocl::getOpenCLAllocator(), u.u->currAllocator);
+#endif
+}
+
+TEST(Core_CudaUMat, op_arithmetic_host_fallback)
+{
+    skipIfNoCudaDevice();
+    NoOpenCL guard;
+
+    Mat a(50, 40, CV_32FC1), b(50, 40, CV_32FC1);
+    randu(a, 0, 10); randu(b, 0, 10);
+
+    UMat ua = makeCudaUMat(), ub = makeCudaUMat();
+    a.copyTo(ua); b.copyTo(ub);
+
+    UMat usum = makeCudaUMat(), udiff = makeCudaUMat(), uprod = makeCudaUMat();
+    cv::add(ua, ub, usum);
+    cv::subtract(ua, ub, udiff);
+    cv::multiply(ua, ub, uprod);
+
+    Mat sum, diff, prod;
+    usum.copyTo(sum); udiff.copyTo(diff); uprod.copyTo(prod);
+    EXPECT_LE(cvtest::norm(sum,  a + b,       NORM_INF), 1e-5);
+    EXPECT_LE(cvtest::norm(diff, a - b,       NORM_INF), 1e-5);
+    EXPECT_LE(cvtest::norm(prod, a.mul(b),    NORM_INF), 1e-4);
+}
+
+TEST(Core_CudaUMat, op_convertTo_host_fallback)
+{
+    skipIfNoCudaDevice();
+    NoOpenCL guard;
+
+    Mat a(30, 45, CV_32FC1);
+    randu(a, 0, 100);
+
+    UMat ua = makeCudaUMat();
+    a.copyTo(ua);
+    UMat u8 = makeCudaUMat();
+    ua.convertTo(u8, CV_8U, 2.0, 1.0);
+
+    Mat got, expected;
+    u8.copyTo(got);
+    a.convertTo(expected, CV_8U, 2.0, 1.0);
+    EXPECT_EQ(0.0, cvtest::norm(got, expected, NORM_INF));
+}
+
+TEST(Core_CudaUMat, op_addWeighted_host_fallback)
+{
+    skipIfNoCudaDevice();
+    NoOpenCL guard;
+
+    Mat a(40, 40, CV_32FC1), b(40, 40, CV_32FC1);
+    randu(a, 0, 5); randu(b, 0, 5);
+
+    UMat ua = makeCudaUMat(), ub = makeCudaUMat();
+    a.copyTo(ua); b.copyTo(ub);
+
+    UMat ur = makeCudaUMat();
+    cv::addWeighted(ua, 2.0, ub, 1.0, 0.0, ur);
+
+    Mat got, expected;
+    ur.copyTo(got);
+    cv::addWeighted(a, 2.0, b, 1.0, 0.0, expected);
+    EXPECT_LE(cvtest::norm(got, expected, NORM_INF), 1e-4);
+}
+
+}} // namespace
+
+#endif // HAVE_CUDA