CUDA Programming Interface

Parallel Algorithms

David Mayerich

Scalable Tissue Imaging and Modeling (STIM) Laboratory

Department of Electrical and Computer Engineering

Cullen College of Engineering

University of Houston

David Mayerich

STIM Laboratory, University of Houston

CUDA API and Language

Application Programming Interface (API)

C/C++ functions that query, initialize, and execute GPU features
Allocating space in device memory, copying data, getting device information
Setting up special types of memory on the GPU

David Mayerich

STIM Laboratory, University of Houston

CUDA Language

Implement code executed on CUDA devices
ANSI C is a subset of CUDA
Additional CUDA keywords:
specify regions of data-parallel code
describe how parallel regions are executed
define how resources are allocated

CUDA Device "Marketecture"

David Mayerich

STIM Laboratory, University of Houston

\cdots

global memory

streaming

multi-

processor

(SM)

streaming multiprocessor (SM)

instruction

shared memory

\cdots

CUDA Kernels

A kernel is a block of code executed in parallel on a CUDA device
Similar to a callback function for an std::thread

David Mayerich

STIM Laboratory, University of Houston

__global__ void mmult(float* A, float* B, float* C, int N) {
  int row = blockIdx.y * blockDim.y + threadIdx.y;
  int col = blockIdx.x * blockDim.x + threadIdx.x;
  
  if(row < N && col < N){
    float sum = 0.0f;
    for(int k = 0; k < N; k++) {
      sum += A[row * N + k] * B[k * N + col];
    }
  }
}

CUDA Threads (vs. `std::threads`)

David Mayerich

STIM Laboratory, University of Houston

Serial Code
(CPU)

Data Parallel
Code

Serial Code

Data Parallel
Code

main()

std::thread[N]

join()

std::thread[M]

join()

main()

launch GPU kernel

synchronize

launch GPU kernel

CUDA Device Properties

Get the number of devices (GPUs) available
- a host (CPU) can have several CUDA devices:
  cudaError_t cudaGetDeviceCount(int* n)
- the programmer can control what devices executes code:
  cudaError_t cudaGetDevice(int* d) cudaError_t cudaSetDevice(int d)
Check for compatibility
- test for compute capability - determine what features the device has
- test for available memory - how much data can be copied to the device?
  cudaError_t cudaGetDeviceProperties(cudaDevProp* prop, int dev);

David Mayerich

STIM Laboratory, University of Houston

CUDA Device Properties

global memory defines how large of a problem you can solve with a single kernel
compute capability specifies the set of features available on the device

David Mayerich

STIM Laboratory, University of Houston

struct cudaDeviceProp {
  char name[256];
  size_t totalGlobalMem;
  // ...
  int major;
  int minor;
  // ...
  size_t sharedMemPerMultiprocessor;
  size_t totalGlobalMem;
  size_t totalConstMem;
  size_t warpSize;
  int ECCEnabled;
  // ...
}

compute	date	features
1.1	2006	global memory atomics
1.3	2008	64-bit floating point
2.x	2010	L1/L2 caches, unified address spaces
5.3	2014	16-bit floating point
6.1	2016	64-bit and 16-bit atomics
7.0	2017	tensor cores
8.9	2022	ray tracing cores

Error Checking

All CUDA API functions return a cudaError_t value

David Mayerich

STIM Laboratory, University of Houston

cudaError_t error;
error = cudaMalloc(gpuPtr, N * sizeof(float));
if( error != cudaSuccess) {
  throw std::runtime_error("Error allocating memory.");
}

cudaError_t =
cudaSuccess // operation successful cudaErrorMemoryAllocation // not enough memory cudaErrorInvalidValue // a parameter is wrong cudaErrorInvalidDevicePointer // bad GPU pointer

Error checking example:

static void HandleError( cudaError_t err, const char* file, int line) {
  if(err != cudaSuccess)
    std::cout<<cudaGetErrorString(err)<<" in "<< file <<" at line "<< line;
}
#define HANDLE_ERROR( err ) (HandleError( err, __FILE__, __LINE__))

Helper macro:

CUDA API: Basic Framework

Query for compatible devices
Allocate CUDA device memory
Copy data from the host
Launch the kernel
Copy data back to the host

David Mayerich

STIM Laboratory, University of Houston

int num;
cudaGetDeviceCount(&num);
int device_id = 0; 		// choose a device
cudaDevProp* prop;
cudaGetDeviceProperties(&prop, device_id);

float* gpu_array;
size_t bytes = sizeof(float) * N;
cudaMalloc(&gpu_array, bytes);

float* array;
// fill array;
cudaMemcpy(gpu_array, array, bytes, 
		   cudaMemcpyHostToDevice);

float* result;
cudaMemcpy(result, gpu_result, bytes, 
		   cudaMemcpyDeviceToHost);

Example (without Error Checking)

David Mayerich

STIM Laboratory, University of Houston

void main(int argc, char* argv[]) {
  int N = 30000;					// allocate an array in main memory
  float* x = new float[N];
  size_t array_size = N * sizeof(float);

  // load array values (ex. from a file)
  
  float* gpu_x;						// create a pointer to device memory
  // allocate memory on the CUDA device
  cudaMalloc(&gpu_x, array_size);
  // copy the array from main memory (host) to the device
  cudaMemcpy(gpu_x, x, array_size, cudaMemcpyHostToDevice);
  
  // run device code (launch kernel)
  
  // copy data from the device to main memory (host)
  cudaMemcpy(x, gpu_x, array_size, cudaMemcpyDeviceToHost);
}

Example (with Error Checking)

David Mayerich

STIM Laboratory, University of Houston

void main(int argc, char* argv[]) {
  int N = 30000;					// allocate an array in main memory
  float* x = new float[N];
  size_t array_size = N * sizeof(float);
  
  // load array values (ex. from a file)
  
  float* gpu_x;						// create a pointer to device memory
  // allocate memory on the CUDA device
  HANDLE_ERROR( cudaMalloc(&gpu_x, array_size));
  // copy the array from main memory (host) to the device
  HANDLE_ERROR( cudaMemcpy(gpu_x, x, array_size, cudaMemcpyHostToDevice));
  
  // run device code (launch kernel)
  
  // copy data from the device to main memory (host)
  HANDLE_ERROR( cudaMemcpy(x, gpu_x, array_size, cudaMemcpyDeviceToHost));
}

CUDA Compiler (`nvcc`)

nvcc strips CUDA code from C/C++
CUDA code is converted to device code
- parallel thread execution (PTX) code is compiled just-in-time for the target device
- runtime compiler converts PTX into device binary code that can be executed directly on the CUDA devices
C/C++ code is passed to the native compiler: gcc, g++, Visual Studio, etc.
- C/C++ code is compiled and linked normally
- object files are linked to CUDA libraries that handle device execution
- This operation is usually transparent to the programmer
- PTX can be useful for profiling and optimization

David Mayerich

STIM Laboratory, University of Houston

CUDA Compiler (`nvcc`)

David Mayerich

STIM Laboratory, University of Houston

standard C/C++ and CUDA code

nvcc compiler

gcc, g++, VS, etc.

CPU/OS

device compiler

void cpu_sum_vec() {

char* a = malloc(N);

memset(a, 128, N);

...

}

__device__ sum_vec(char* a, N) {

i = threadIdx.x;

c[i] = a[i] + b[i];

...

}

...

char* a = malloc(N);

memset(a, 128, N);

...

ld.s32 r1, [j];

ld.s32 r2, [k];

add.s32 r3, r1, r2;

st.s32 [i], r3;

1010100101001010100

0101010010100101010

0101001001010010101

1110010010010101110

1010100101001010100

0101010010100101010

0101001001010010101

1110010010010101110

C/C++

CPU binary

device code

(PTX)

JIT

device binary

CUDA Thread Resources

Threads are executed in parallel on the GPU
Tesla P100 has 56 SMs, each with 64 cores = 3584 total CUDA cores
- at least this many instructions can be executed simultaneously
- additional threads can be accommodated with hyperthreading
Some resources are shared:
- global memory is shared across the entire device (all threads)
- shared memory is shared across a single SM
- registers are local to a thread
Proximity = speed:
- global memory (200 - 400 clock cycle latency)
- shared memory (20 - 40 clock cycle latency)
- registers are fast (1 - 11 clock cycles depending on the instruction)

David Mayerich

STIM Laboratory, University of Houston

Declaring Kernel Functions

CUDA contains three additional keywords for function declarations
Declarations specify how the code will be compiled by nvcc

David Mayerich

STIM Laboratory, University of Houston

__host__ // called and executed on the host

__device__ // called and executed on a device

__global__ // called on the host, executed on a device

Keywords start the declaration:

__host__ int abs(int v) {
  if(v < 0) return -v;
  return v;
}

__device__ int abs(int v) {
  if(v < 0) return -v;
  return v;
}

__host__ __device__ int abs(int v) {
  if(v < 0) return -v;
  return v;
}

#define CUDA_CALLABLE __host__ __device__
CUDA_CALLABLE int abs(int v) {
  if(v < 0) return -v;
  return v;
}

Both can be used (compiles 2 functions)

host Functions

Mostly ignored by nvcc
- kernels can be launched from host functions
- launches are stripped from C/C++ and replaced with CUDA launch instructions
__host__ functions can only be called by other __host__ functions
Functions without CUDA declarations are assumed to be __host__ functions

David Mayerich

STIM Laboratory, University of Houston

int abs(int v) {
  if(v < 0) return -v;
  return v;
}

__host__ int abs(int v) {
  if(v < 0) return -v;
  return v;
}

is identical to

device Functions

Compiled completely by nvcc (not passed to the host compiler)
Only executable on a device
Only __device__ functions can be called by __device__ functions
Once you're executing on the device, the only way off is to return

David Mayerich

STIM Laboratory, University of Houston

__host__ __device__ int abs(int v) {
  if(v < 0) return -v;
  return v;
}

#define CUDA_CALLABLE __host__ __device__
CUDA_CALLABLE int abs(int v) {
  if(v < 0) return -v;
  return v;
}

Functions can be declared as both __host__ and __device__
Two versions of the function are created: one in PTC (by nvcc) and the other by the native compiler

global Functions

Used to specify kernels (functions that execute on the GPU)
Kernels are called by __host__ functions
- special parameters are specified to define how the kernel is launched
- once launched, the kernel runs on the device
- execution continues on the host function that launched the kernel
__global__ functions are device functions called from __host__ functions
__global__ functions must have a void return type

David Mayerich

STIM Laboratory, University of Houston

CUDA Programming Interface

Parallel Algorithms

David Mayerich

CUDA API and Language

CUDA Device "Marketecture"

CUDA Kernels

CUDA Threads (vs. std::threads)

CUDA Device Properties

CUDA Device Properties

Error Checking

CUDA API: Basic Framework

Example (without Error Checking)

Example (with Error Checking)

CUDA Compiler (nvcc)

CUDA Compiler (nvcc)

CUDA Thread Resources

Declaring Kernel Functions

__host__ Functions

__device__ Functions

__global__ Functions

CUDA Threads (vs. `std::threads`)

CUDA Compiler (`nvcc`)

CUDA Compiler (`nvcc`)

host Functions

device Functions

global Functions