I’d like to show how to use an HPC part written in C++ with CUDA from Python code. So, every heavy part can be done on the GPU with CUDA, while all gluing tasks (with beautiful matplotlib plots) are done on the CPU with Python.
We will use a shared object compiled from C++ CUDA code in Python. The only uncertain part here is the conversion of types from «high-level» Python ones to «low-level» C++ ones. We will write an application for the parallel calculation of the elementwise sum of two arrays.
First, the CUDA code. The CUDA kernel cuda_sum_kernel does the job on the GPU, while the wrapper cuda_sum prepares arrays for the GPU and frees memory after the calculation is done. Note the extern "C" line. It is important for the correct function name in the compiled shared object later.
1#include <cuda.h>
2#include <cuda_runtime_api.h>
3
4__global__ void cuda_sum_kernel(float *a, float *b, float *c, size_t size)
5{
6 size_t idx = blockIdx.x * blockDim.x + threadIdx.x;
7 if (idx >= size) {
8 return;
9 }
10
11 c[idx] = a[idx] + b[idx];
12}
13
14extern "C" {
15void cuda_sum(float *a, float *b, float *c, size_t size)
16{
17 float *d_a, *d_b, *d_c;
18
19 cudaMalloc((void **)&d_a, size * sizeof(float));
20 cudaMalloc((void **)&d_b, size * sizeof(float));
21 cudaMalloc((void **)&d_c, size * sizeof(float));
22
23 cudaMemcpy(d_a, a, size * sizeof(float), cudaMemcpyHostToDevice);
24 cudaMemcpy(d_b, b, size * sizeof(float), cudaMemcpyHostToDevice);
25
26 cuda_sum_kernel <<< ceil(size / 256.0), 256 >>> (d_a, d_b, d_c, size);
27
28 cudaMemcpy(c, d_c, size * sizeof(float), cudaMemcpyDeviceToHost);
29
30 cudaFree(d_a);
31 cudaFree(d_b);
32 cudaFree(d_c);
33}
34}
Compile it to a *.so file with the nvcc compiler:
1/usr/local/cuda/bin/nvcc -Xcompiler -fPIC -shared -o cuda_sum.so cuda_sum.cu
The last part is to use the cuda_sum function from the created cuda_sum.so file in a Python script. Example (with comments):
1import numpy as np
2import ctypes
3from ctypes import *
4
5# extract cuda_sum function pointer in the shared object cuda_sum.so
6def get_cuda_sum():
7 dll = ctypes.CDLL('./cuda_sum.so', mode=ctypes.RTLD_GLOBAL)
8 func = dll.cuda_sum
9 func.argtypes = [POINTER(c_float), POINTER(c_float), POINTER(c_float), c_size_t]
10 return func
11
12# create __cuda_sum function with get_cuda_sum()
13__cuda_sum = get_cuda_sum()
14
15# convenient python wrapper for __cuda_sum
16# it does all job with types convertation
17# from python ones to C++ ones
18def cuda_sum(a, b, c, size):
19 a_p = a.ctypes.data_as(POINTER(c_float))
20 b_p = b.ctypes.data_as(POINTER(c_float))
21 c_p = c.ctypes.data_as(POINTER(c_float))
22
23 __cuda_sum(a_p, b_p, c_p, size)
24
25# testing, sum of two arrays of ones and output head part of resulting array
26if __name__ == '__main__':
27 size=int(1024*1024)
28
29 a = np.ones(size).astype('float32')
30 b = np.ones(size).astype('float32')
31 c = np.zeros(size).astype('float32')
32
33 cuda_sum(a, b, c, size)
34
35 print c[:10]
Now you can design the code above into a full-featured Python module.