Question

from numbapro import cuda, vectorize, guvectorize, check_cuda
from numbapro import void, uint8 , uint32, uint64, int32, int64, float32, float64, f8
import numpy as np

check_cuda()

------------------------------libraries detection-------------------------------
Finding cublas
    located at /Users/cliburn/anaconda/lib/libcublas.6.0.dylib
    trying to open library...       ok
Finding cusparse
    located at /Users/cliburn/anaconda/lib/libcusparse.6.0.dylib
    trying to open library...       ok
Finding cufft
    located at /Users/cliburn/anaconda/lib/libcufft.6.0.dylib
    trying to open library...       ok
Finding curand
    located at /Users/cliburn/anaconda/lib/libcurand.6.0.dylib
    trying to open library...       ok
Finding nvvm
    located at /Users/cliburn/anaconda/lib/libnvvm.2.0.0.dylib
    trying to open library...       ok
    finding libdevice for compute_20...     ok
    finding libdevice for compute_30...     ok
    finding libdevice for compute_35...     ok
-------------------------------hardware detection-------------------------------
Found 1 CUDA devices
id 0         GeForce GTX 760                              [SUPPORTED]
                      compute capability: 3.0
                           pci device id: 0
                              pci bus id: 1
Summary:
    1/1 devices are supported
PASSED

True

Let’s start by doing vector addition on the GPU with a kernel function. This requires several steps:

1. Define the kernel function(s) (code to be run on parallel on the GPU)
   1. In simplest model, one kernel is executed at a time and then control returns to CPU
   2. Many threads execute one kernel
2. Allocate space on the CPU for the vectors to be added and the solution vector
3. Copy the vectors onto the GPU
4. Run the kernel with grid and blcok dimensions
5. Copy the solution vector back to the CPU

Image(url='https://code.msdn.microsoft.com/vstudio/site/view/file/95904/1/Grid-2.png')

Execution rules:

• All threads in a grid execute the same kernel function
• A grid is organized as a 2D array of blocks
• All blocks in a grid have the same dimension
• Total size of a block is limited to 512 or 1024 threads

Definitions:

• gridDim: This variable contains the dimensions of the grid (gridDim.x and gridDim.y)
• blockIdx: This variable contains the block index within the grid
• blockDim: This variable and contains the dimensions of the block (blockDim.x, blockDim.y and blockDim.z)
• threadIdx: This variable contains the thread index within the block.

How do we find out the unique global thread identity?

To execute kernels in parallel with CUDA, we launch a grid of blocks of threads, specifying the number of blocks per grid ( bpg ) and threads per block ( tpb ). The total number of threads launched will be the product of bpg \(\times\) tpb . This can be in the millions.

Now, in order to decide what thread is doing what, we need to find its gloabl ID. This is basically just finding an offset given a 2D grid of 3D blocks of 3D threads, but can get very confusing.

1D grid of 1D blocks

bx = cuda.blockIdx.x
bw = cuda.blockDim.x
tx = cuda.threadIdx.x
i = tx + bx * bw

2D grid of 2D blocsk

tx = cuda.threadIdx.x
ty = cuda.threadIdx.y
bx = cuda.blockIdx.x
by = cuda.blockIdx.y
bw = cuda.blockDim.x
bh = cuda.blockDim.y
i = tx + bx * bw
j = ty + by * bh

3D grid of 3D blocks

tx = cuda.threadIdx.x
ty = cuda.threadIdx.y
tz = cuda.threadIdx.z
bx = cuda.blockIdx.x
by = cuda.blockIdx.y
bz = cuda.blockIdx.y
bw = cuda.blockDim.x
bh = cuda.blockDim.y
bd = cuda.blockDim.z
i = tx + bx * bw
j = ty + by * bh
k = tz + bz * hd

More exotic combinations - e.g. 3D grid of 2D blockss are also possible but uncommon. If you do have a problem that masp to one of these geometrires, see this cheatshet for calculating the global thread index.

Fortunately, these \(1 \times 1\), \(2 \times 2\) and \(3 \times 3\) patterns are so common that theere is a shorthand macro proivded in CUDA Python using the grid macro.