可扩展的编程模型使得已编译好的 CUDA 程序能够在任意核心的 GPU 上执行。

内核函数

下面这段代码通过语法 <<<...>>> 指定使用 N 个线程去执行 VecAdd 函数，每个执行内核函数的线程拥有一个独一无二的线程 ID，可以通过内置的 threadIdx 变量在内核函数中访问。限定符 __global__ 声明某函数为内核函数，在设备上执行，只能从主机端调用。

__global__ void VecAdd(float* A, float* B, float* C) {
    int i = threadIdx.x; C[i] = A[i] + B[i];
}

int main() {
    ...
    VecAdd<<<1, N>>>(A, B, C);
    ...
}

线程

下面这段代码指定使用 1 个含有 N*N 个 线程（二维网格形状）的线程块执行 MatAdd 函数。另外，一个块内的线程数目是有限的，在目前的 GPU 中，一个线程块最多有 1024 个线程。

__global__ void MatAdd(float A[N][N], float B[N][N], float C[N][N]) {
    int i = threadIdx.x;
    int j = threadIdx.y;
    C[i][j] = A[i][j] + B[i][j];
}

int main() {
    ...
    int numBlocks = 1;
    dim3 threadsPerBlock(N, N);
    MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
    ...
}

一个内核函数可被多个同样大小的线程块执行。下面这段代码指定使用 N/threadsPerBlock.x*N/threadsPerBlock.y 个含有 16*16 个线程的线程块执行 MatAdd 函数。

__global__ void MatAdd(float A[N][N], float B[N][N], float C[N][N]) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    int j = blockIdx.y * blockDim.y + threadIdx.y;
    if (i < N && j < N)
        C[i][j] = A[i][j] + B[i][j];
}

int main() {
    ...
    dim3 threadsPerBlock(16, 16);
    dim3 numBlocks(N / threadsPerBlock.x, N / threadsPerBlock.y);
    MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
    ...
}

参考

CUDA C++ Programming Guide
《CUDA C Programming Guide》(《CUDA C 编程指南》)导读
《CUDA 并行程序设计：GPU 编程指南》