Preface
The standard CUDA programming model can be broadly abstracted into a basic flow like “Prepare&Launch => Execute => Sync => ReadBack”. In an ideal case, as long as the workload is well-suited to data parallelism, there is no severe imbalance in partitioning, and the data granularity is not too small, the GPU’s multi-core compute capability can be utilized effectively. However, if the compute capability of the underlying device is very high, and the time spent per computation is almost on the same order as a kernel launch (~20us), then this pattern of frequent interaction with the host introduces relatively high overhead. In a single-worker-thread scenario, it may fail to fully utilize the underlying hardware. This runtime overhead will become increasingly apparent as heterogeneous compute devices continue to improve.
Persistent Thread (PT for short below) is an important CUDA optimization technique that can substantially reduce GPU “kernel launch latency” and the extra overhead introduced by Host-Device communication. However, due to programming model constraints stemming from GPU hardware design characteristics, this technique has not become widely adopted (see the survey paper [2] for details). This article analyzes the working principles, application scenarios, code implementation, pros and cons, and the minimal feature subset required to support this technique.
In addition, CUDA officially provides a “Dynamic Parallelism” mechanism to better support special cases like “Data-Dependent Parallelism”. The problems it addresses overlap slightly with PT. For details, see this simplified slides and the corresponding blog.
Persistent Thread
TL; DR: Essentially, the core idea of the PT technique is to keep the same kernel code running on the GPU device, and use UVA[1] or zero-copy memory to communicate with the GPU at extremely low overhead, offloading workloads to a kernel on the compute device that continuously polls for tasks. In other words, it treats the external compute device as a CPU-visible thread pool. In contrast, loading programs via kernel launch (normal Host-Device interaction) requires going through a full path from the upper runtime down to the PCIE driver, which is much more expensive.
Basic Application Scenarios
CPU-GPU Synchronization
Load Balancing / Irregular Parallelism
Producer-Consumer Locality
Global Synchronization
Code Implementation
Below, we analyze specific application scenarios of the Persistent Kernel programming model using concrete code examples.
Device Communication
Consider the following code:
// test.cu
// nvcc test.cu -o test && ./test
#include <stdio.h>
#include <stdlib.h>
#include <time.h>
#define TIME_INC 100000000
#define INC_CNT 10
#define cudaCheckErrors(msg) \
do { \
cudaError_t __err = cudaGetLastError(); \
if (__err != cudaSuccess) { \
fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \
msg, cudaGetErrorString(__err), \
__FILE__, __LINE__); \
fprintf(stderr, "*** FAILED - ABORTING\n"); \
exit(1); \
} \
} while (0)
__global__ void test_kernel(volatile int *data) {
unsigned long time;
for (int i = 0; i < INC_CNT; i++) {
atomicAdd((int *)data, 1);
__threadfence_system();
time = clock64();
while ((clock64() - time) < TIME_INC) {};
}
printf("demo kernel exited\n");
}
int main() {
volatile int *d_data, *h_data;
cudaSetDeviceFlags(cudaDeviceMapHost);
cudaCheckErrors("cudaSetDeviceFlags error");
cudaHostAlloc((void **)&h_data, sizeof(int), cudaHostAllocMapped);
cudaCheckErrors("cudaHostAlloc error");
cudaHostGetDevicePointer((int **)&d_data, (int *)h_data, 0);
cudaCheckErrors("cudaHostGetDevicePointer error");
*h_data = 0;
printf("Kernel starting\n");
test_kernel<<<1, 1>>>(d_data);
cudaCheckErrors("kernel failed");
int value = 0;
do {
int old_val = *h_data;
if (old_val > value) {
printf("h_data = %d\n", old_val);
value = old_val;
}
} while (value < (INC_CNT - 1));
cudaDeviceSynchronize();
cudaCheckErrors("kernel failed");
return 0;
}
The runtime output of the above code is:
$ ./test
Kernel starting
h_data = 1
h_data = 2
h_data = 3
h_data = 4
h_data = 5
h_data = 6
h_data = 7
h_data = 8
h_data = 9
progress check finished
Now consider a slightly more complex Host-Device synchronization. Here we use a progress variable to notify the host of kernel execution progress in real time:
__global__ void
matmult(float *a, float *b, float *c,
unsigned int rowA, unsigned int colA,
unsigned int colB, volatile int *progress) {
unsigned int row = threadIdx.x + blockDim.x * blockIdx.x;
unsigned int col = threadIdx.y + blockDim.y * blockIdx.y;
if ((row < rowA) && (col < colB)) {
float temp = 0.0f;
for (unsigned int k = 0; k < colA; k++) {
temp += a[(row * colA) + k] * b[(k * colB) + col];
}
c[(row * colB) + col] = temp;
if (!(threadIdx.x || threadIdx.y)) {
atomicAdd((int *)progress, 1);
__threadfence_system();
}
}
}
int main() {
int value = 0;
volatile int *d_data, *h_data;
cudaSetDeviceFlags(cudaDeviceMapHost);
cudaCheckErrors("cudaSetDeviceFlags error");
cudaHostAlloc((void **)&h_data, sizeof(int), cudaHostAllocMapped);
cudaCheckErrors("cudaHostAlloc error");
cudaHostGetDevicePointer((int **)&d_data, (int *)h_data, 0);
cudaCheckErrors("cudaHostGetDevicePointer error");
float *h_c, *d_a, *d_b, *d_c;
h_c = (float *)malloc(MAT_DIMX * MAT_DIMY * sizeof(float));
if (h_c == NULL) {
printf("malloc fail\n");
return 1;
}
cudaMalloc((void **)&d_a, MAT_DIMX * MAT_DIMY * sizeof(float));
cudaCheckErrors("cudaMalloc a fail");
cudaMalloc((void **)&d_b, MAT_DIMX * MAT_DIMY * sizeof(float));
cudaCheckErrors("cudaMalloc b fail");
cudaMalloc((void **)&d_c, MAT_DIMX * MAT_DIMY * sizeof(float));
cudaCheckErrors("cudaMalloc c fail");
for (int i = 0; i < MAT_DIMX * MAT_DIMY; i++) { h_c[i] = rand() / (float)RAND_MAX; }
cudaMemcpy(d_a, h_c, MAT_DIMX * MAT_DIMY * sizeof(float), cudaMemcpyHostToDevice);
cudaCheckErrors("cudaMemcpy a fail");
cudaMemcpy(d_b, h_c, MAT_DIMX * MAT_DIMY * sizeof(float), cudaMemcpyHostToDevice);
cudaCheckErrors("cudaMemcpy b fail");
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
*h_data = 0;
dim3 block(16, 16);
dim3 grid(((MAT_DIMX + block.x - 1) / block.x), ((MAT_DIMY + block.y - 1) / block.y));
printf("matrix multiply kernel starting\n");
cudaEventRecord(start);
matmult<<<grid,block>>>(d_a, d_b, d_c, MAT_DIMY, MAT_DIMX, MAT_DIMX, d_data);
cudaEventRecord(stop);
unsigned int num_blocks = grid.x * grid.y;
float my_progress = 0.0f;
value = 0;
printf("Progress:\n");
do {
int value1 = *h_data;
float kern_progress = (float)value1 / (float)num_blocks;
if ((kern_progress - my_progress) > 0.1f) {
printf("percent complete = %2.1f\n", (kern_progress * 100));
my_progress = kern_progress;
}
} while (my_progress < 0.9f);
printf("\n");
cudaEventSynchronize(stop);
cudaCheckErrors("event sync fail");
float et;
cudaEventElapsedTime(&et, start, stop);
cudaCheckErrors("event elapsed time fail");
cudaDeviceSynchronize();
cudaCheckErrors("mat mult kernel fail");
printf("matrix multiply finished. elapsed time = %f milliseconds\n", et);
return 0;
}
The execution output of the above code is:
$ ./test
matrix multiply kernel starting
Progress:
percent complete = 10.0
percent complete = 20.0
percent complete = 30.0
percent complete = 40.0
percent complete = 50.0
percent complete = 60.0
percent complete = 70.0
percent complete = 80.0
percent complete = 90.0
matrix multiply finished. elapsed time = 6264.967285 milliseconds
In fact, on newer CUDA versions, the pointer allocated by cudaHostAlloc can be passed directly to the kernel. This pointer can be dereferenced by the GPU device via the UVM mechanism, and there is no need to call cudaHostGetDevicePointer to obtain a device-visible pointer. The extra step here is only to illustrate the concept from earlier programming models.
Task Dispatch
Ping-Pong Buffer
// test.cu
// nvcc test.cu -o test && ./test
#include <stdio.h>
#define TOTAL_ITERS 1000
#define DSIZE 65536
#define nTPB 256
#define BUFFER_EMPTY 0
#define BUFFER_FULL 1
#define cudaCheckErrors(msg) \
do { \
cudaError_t __err = cudaGetLastError(); \
if (__err != cudaSuccess) { \
fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \
msg, cudaGetErrorString(__err), \
__FILE__, __LINE__); \
fprintf(stderr, "*** FAILED - ABORTING\n"); \
exit(1); \
} \
} while (0)
__device__ void do_compute(int *buf, int idx, int data) {
buf[idx] = data;
}
__global__ void test_kernel(int *buffer1, int *buffer2,
volatile int *buffer1_ready,
volatile int *buffer2_ready,
const int buffersize) {
static volatile int completed_iter = 0;
static volatile int blkcnt1 = 0;
static volatile int blkcnt2 = 0;
// assumption of persistent block-limited kernel launch
int idx = threadIdx.x + blockDim.x * blockIdx.x;
// persistent until TOTAL_ITERS complete
for (int current_iter = 0; current_iter < TOTAL_ITERS; current_iter++) {
int *buf = (current_iter & 1) ? buffer2 : buffer1; // ping pong between buffers
volatile int *bufrdy = (current_iter & 1) ? (buffer2_ready) : (buffer1_ready);
volatile int *blkcnt = (current_iter & 1) ? (&blkcnt2) : (&blkcnt1);
int my_idx = idx;
while (current_iter - completed_iter > 1); // don't overrun buffers on device
while (*bufrdy != BUFFER_EMPTY); // wait for buffer to be consumed
while (my_idx < buffersize) { // perform the "work"
do_compute(buf, my_idx, current_iter);
my_idx += gridDim.x * blockDim.x; // grid-striding loop
}
__syncthreads(); // wait for my block to finish
__threadfence(); // make sure global buffer writes are "visible"
// now do global sync
if (!threadIdx.x) { atomicAdd((int *)blkcnt, 1); } // mark my block done
if (!idx) { // am I the master block/thread?
while (*blkcnt < gridDim.x); // wait for all blocks to finish
*blkcnt = 0;
*bufrdy = BUFFER_FULL; // indicate that buffer is ready
__threadfence_system(); // push it out to mapped memory
completed_iter++;
}
}
}
int validate(const int *data, const int dsize, const int val) {
for (int i = 0; i < dsize; i++) {
if (data[i] != val) {
printf("mismatch at %d, was: %d, should be: %d\n", i, data[i], val);
return 0;
}
}
return 1;
}
int main() {
int *h_buf1, *d_buf1, *h_buf2, *d_buf2;
volatile int *m_bufrdy1, *m_bufrdy2;
// buffer and "signal" setup
cudaHostAlloc(&h_buf1, DSIZE * sizeof(int), cudaHostAllocDefault);
cudaHostAlloc(&h_buf2, DSIZE * sizeof(int), cudaHostAllocDefault);
cudaHostAlloc(&m_bufrdy1, sizeof(int), cudaHostAllocMapped);
cudaHostAlloc(&m_bufrdy2, sizeof(int), cudaHostAllocMapped);
cudaCheckErrors("cudaHostAlloc fail");
cudaMalloc(&d_buf1, DSIZE * sizeof(int));
cudaMalloc(&d_buf2, DSIZE * sizeof(int));
cudaCheckErrors("cudaMalloc fail");
cudaStream_t streamk, streamc;
cudaStreamCreate(&streamk);
cudaStreamCreate(&streamc);
cudaCheckErrors("cudaStreamCreate fail");
*m_bufrdy1 = BUFFER_EMPTY;
*m_bufrdy2 = BUFFER_EMPTY;
cudaMemset(d_buf1, 0xFF, DSIZE * sizeof(int));
cudaMemset(d_buf2, 0xFF, DSIZE * sizeof(int));
cudaCheckErrors("cudaMemset fail");
// inefficient crutch for choosing number of blocks
int nblock = 0;
cudaDeviceGetAttribute(&nblock, cudaDevAttrMultiProcessorCount, 0);
cudaCheckErrors("cudaDeviceGetAttribute fail");
test_kernel<<<nblock, nTPB, 0, streamk>>>(d_buf1, d_buf2, m_bufrdy1, m_bufrdy2, DSIZE);
cudaCheckErrors("kernel launch fail");
volatile int *bufrdy;
int *hbuf, *dbuf;
for (int i = 0; i < TOTAL_ITERS; i++) {
// ping pong on the host side
if (i & 1) {
bufrdy = m_bufrdy2;
hbuf = h_buf2;
dbuf = d_buf2;
} else {
bufrdy = m_bufrdy1;
hbuf = h_buf1;
dbuf = d_buf1;
}
while ((*bufrdy) != BUFFER_FULL);
cudaMemcpyAsync(hbuf, dbuf, DSIZE * sizeof(int), cudaMemcpyDeviceToHost, streamc);
cudaStreamSynchronize(streamc);
cudaCheckErrors("cudaMemcpyAsync fail");
*bufrdy = BUFFER_EMPTY; // release buffer back to device
if (!validate(hbuf, DSIZE, i)) {
printf("validation failure at iter %d\n", i);
exit(1);
}
}
printf("Completed %d iterations successfully\n", TOTAL_ITERS);
return 0;
}
There are several noteworthy points in the above code:
Pointer passing
We directly pass the pointer allocated by cudaHostAlloc to the kernel, because this memory is visible to both the host and the GPU and atomic operations can be performed on it directly; no additional special handling is required.
Stream concurrency
We create two different streams for kernel execution and the subsequent data readback memcpy, ensuring they do not block each other.
Choosing the number of blocks
When launching the kernel, we start the same number of blocks as the number of SMs on the device, to ensure these blocks are scheduled and started simultaneously; otherwise, global synchronization may break. In fact, this strategy is the safest but not optimal, because multiple blocks can be scheduled onto one SM, and this upper limit is not 1. How many blocks an SM can handle depends heavily on resource usage, so this number is closely related to the kernel implementation. In addition, each block launches a fixed 256 threads, but this has limited impact because threads are time-sliced anyway, while blocks are not.
Global synchronization
Within each thread, after processing its own data, we first perform a per-block synchronization, and then a global synchronization across all blocks to ensure that all data for the current task is processed before moving on to the next task. Due to limitations of the CUDA programming model, this can likewise only be implemented using atomic variables. If we launch more blocks than the device can schedule concurrently, then not all blocks will be running during synchronization, and the synchronization will fail and result in an infinite loop.
Synchronization overhead
If we set the workload to empty and measure only synchronization overhead, we will observe some interesting phenomena. Overall, synchronization time is positively correlated with the number of threads and the number of blocks. When only one block is launched and the number of threads <= 64, synchronization overhead is minimal, which matches our understanding of SMs: each SM has two warp dispatchers, which can support exactly 64 physical threads; any more becomes serialized. On the other hand, the number of blocks and synchronization overhead are not strictly proportional.
Producer-Consumer Queue
(To be completed)
Some Thoughts
The first-generation Intel Xeon Phi compute accelerator (KNC) rather strangely adopted a GPU-like working mode: using the host CPU to control the execution flow, while the many fully functional (whether performance matches a CPU is another question) x86 cores on the card only participate in computation. The result was that product generation essentially died. PS: Tianhe-2 was full of these (Xeon Phi 31S1P); in hindsight, it is fair to say it was a costly mistake...
The second-generation KNL was much smarter: it was directly a many-core CPU, handling both scheduling and computation itself. When users write programs, they do not need to worry about which code runs on the CPU versus on the accelerator card; even existing MPI programs can be applied seamlessly. It seems to have gained some market acceptance.
So my personal guess is whether heterogeneous computing will ultimately converge toward homogeneity at the programming-model level: even if processor cores are heterogeneous, some high-performance hardware interconnect could minimize the communication overhead between the accelerator device and the logical control unit as much as possible, making it pleasant for developers to use...
References
- Nvidia Blog: Unified Virtual Addressing
- PT survey paper: Persistent-Threads-Style-Programming-Model-for-GPU-Computing, and the corresponding simplified slides。
- Technical document: Improving Real-Time Performance with CUDA Persistent Threads (CuPer) on the Jetson TX2
- Code example: A minimum CUDA persistent thread example
- Stack Overflow: host-device synchronize state via atomic instructions
- Stack Overflow: ping-pong buffer to implement data interaction between kernel and host