Sorry, something went wrong. Reload?
Sorry, we cannot display this file.
Sorry, this file is invalid so it cannot be displayed.
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
import cupy | |
x = cupy.arange(4096 * 1024, dtype=cupy.float32).reshape((4096, 1024)) | |
y = cupy.arange(1024, dtype=cupy.float32) | |
res_cupy = x + y |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
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]; |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
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)) |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
res_add_broadcast = dispatch_add_broadcast(x, y) |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
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) |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
res = da.map_blocks(dispatch_add_broadcast, dx, dy, dtype=cupy.float32) |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
res.compute() | |
array([[0.000000e+00, 2.000000e+00, 4.000000e+00, ..., 2.042000e+03, | |
2.044000e+03, 2.046000e+03], | |
[1.024000e+03, 1.026000e+03, 1.028000e+03, ..., 3.066000e+03, | |
3.068000e+03, 3.070000e+03], | |
[2.048000e+03, 2.050000e+03, 2.052000e+03, ..., 4.090000e+03, | |
4.092000e+03, 4.094000e+03], | |
..., | |
[4.191232e+06, 4.191234e+06, 4.191236e+06, ..., 4.193274e+06, | |
4.193276e+06, 4.193278e+06], |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
<!DOCTYPE html> | |
<html lang="en"> | |
<head> | |
<meta charset="utf-8"> |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
<!DOCTYPE html> | |
<html lang="en"> | |
<head> | |
<meta charset="utf-8"> |