pentschev

import dask.array as da

dx = da.from_array(x, chunks=(1024, 512), asarray=False)
dy = da.from_array(y, chunks=(1, 512), asarray=False)

pentschev

res_add_broadcast = dispatch_add_broadcast(x, y)

pentschev

def dispatch_add_broadcast(x, y):
    block_size = (32, 32)
    grid_size = (x.shape[1] // block_size[1], x.shape[0] // block_size[0])

    z = cupy.empty(x.shape, x.dtype)

    xdim0 = x.strides[0] // x.strides[1]
    zdim0 = z.strides[0] // z.strides[1]

    add_broadcast_kernel(grid_size, block_size, (x, y, z, xdim0, zdim0))

pentschev

add_broadcast_kernel = cupy.RawKernel(
    r'''
    extern "C" __global__
    void add_broadcast_kernel(
        const float* x, const float* y, float* z,
        const int xdim0, const int zdim0)
    {
        int idx0 = blockIdx.x * blockDim.x + threadIdx.x;
        int idx1 = blockIdx.y * blockDim.y + threadIdx.y;
        z[idx1 * zdim0 + idx0] = x[idx1 * xdim0 + idx0] + y[idx0];

pentschev

import cupy

x = cupy.arange(4096 * 1024, dtype=cupy.float32).reshape((4096, 1024))
y = cupy.arange(1024, dtype=cupy.float32)

res_cupy = x + y

pentschev

from dask.distributed import Client
from dask_cuda import LocalCUDACluster
from dask.array.utils import assert_eq

import dask.array as da
import cupy

add_broadcast_kernel = cupy.RawKernel(
    r'''
    extern "C" __global__

pentschev

--- ucx_info_b	2019-08-08 11:01:32.820637475 -0700
+++ ucx_from_source_info_b	2019-08-08 11:02:15.976537113 -0700
@@ -8,6 +8,7 @@
 #define HAVE_CUDA                 1
 #define HAVE_CUDA_H               1
 #define HAVE_CUDA_RUNTIME_H       1
+#define HAVE_DC_EXP               1
 #define HAVE_DECL_ASPRINTF        1
 #define HAVE_DECL_BASENAME        1
 #define HAVE_DECL_CPU_ISSET       1

pentschev

pentschev

pentschev