Gaps in writing kernel features

[akharche]

Study the Writing a CUDA Kernel section and document with examples equivalent features in numba-dppy.

Feature	CUDA	DPPY
Kernel declaration	@cuda.jit	@numba.jit, @dppy.kernel
Kernel invocation	threadsperblock = 32 blockspergrid = (an_array.size + (threadsperblock - 1)) // threadsperblock increment_by_one [blockspergrid, threadsperblock](an_array)	idx=dppy.get_global_id(0)
Choosing the block size
Thread positioning
Thread index	numba.cuda.threadIdx	dppy.get_local_id
Block dimension	numba.cuda.blockDim	dppy.get_local_size
Block index	numba.cuda.blockIdx	dppy.get_group_id
Grid dimension	numba.cuda.gridDim	dppy.get_num_groups
Grid position	numba.cuda.grid(ndim)	???
Grid size	numba.cuda.gridsize(ndim)	???

Missing features:

• Function of getting absolute position of the current work-item #199
• Function of getting the absolute size (or shape) in work-items #200

[akharche]

Example Numba.cuda:

@cuda.jit
def increment_by_one(an_array):
    # Thread id in a 1D block
    tx = cuda.threadIdx.x
    # Block id in a 1D grid
    ty = cuda.blockIdx.x
    # Block width, i.e. number of threads per block
    bw = cuda.blockDim.x
    # Compute flattened index inside the array
    pos = tx + ty * bw
    if pos < an_array.size:  # Check array boundaries
        an_array[pos] += 1

Example Numba-dppy:

@dppy.kernel
def increment_by_one(an_array, result_array):
    i = dppy.get_global_id(0)
    result_array[i] = an_array[i] + 1