Posts

Showing posts with the label reduction

CUDA - Parallel Reduction over one axis

Image
Clash Royale CLAN TAG #URR8PPP CUDA - Parallel Reduction over one axis I am fairly new to CUDA programming and I am trying to write a CUDA-Kernel for parallel reduction over only 1 dimension of a 3-dimensional tensor which is a row-major flattened float array fed into the kernel. In other words, I am trying to rewrite numpy.sum with limited axes combinations of axis=0, axis=1 and axis=2. I have successfully implemented "reduce over axis 0" and "reduce over axis 1" but performance issues for "reduce over axis2" made me post a question here to ask for advice. The kernel is launched with a 1-D grid and 1-D block configuration and it maps each thread into each element of reduced output tensor. So, it should be something like this: Link to image Here is my kernel: __global__ void kernel_reduce_sum_3d_try02( float* g_idata, float* g_odata, int dim0, int dim1, int dim2, int overaxis0, int overaxis1, ...