并行编程实战——CUDA编程的范式

原创已于 2025-09-28 09:03:35 修改 · 694 阅读

20 ·

CC 4.0 BY-SA版权

文章标签：

#CUDA #c++

于 2025-09-27 10:49:10 首次发布

C++ 同时被 2 个专栏收录

455 篇文章

订阅专栏

并行编程

47 篇文章

订阅专栏

一、编程范式及CUDA的编程范式

在前面学习和分析过什么是范式编程，也对其在C++编程中的应用进行了较为详细的说明。同时，在其它的一些文章中也对各种不同的范式编程进行过介绍，如事件驱动、异步等等编程范式。范式本身是一种抽象，是一种在某个范围内可广泛复用的体系。
而CUDA编程的范式就很容易被说明，它就是异构编程（Heterogeneous Programming）或异构计算（Heterogeneous computing）。为什么这样说呢？因为对大多数的开发者来说，无论是在嵌入式还是在计算机上编程，其底层的架构基本一致。基本上都是围绕着一个核心即CPU或类CPU单元进行工作。但CUDA编程不但涉及到CPU编程还涉及到了GPU编程，更重要的是，它还会把二者结合起来，从而达到应用的目的。

二、CUDA编程范式的基础

通过前面的分析和学习，应该已经基本明白了CUDA编程范式的基础和主要特点，CUDA编程的基础就是有机的融合CPU和GPU的编程。其主要特点包括:

有机的处理CPU和GPU的任务分离
支持大规模的并行计算
层次清晰合理的线程分层模型
多级安全的内存管理体系
提供在CPU和GPU间多种数据交换接口
支持NVIDA的SIMT体系

其实从上面来看，开发者需要重点掌握的主要是线程模型、内存模型和数据交换管理。线程模型和内存模型在前面已经详细的分析过，而数据交换会在后面陆续的展开分析说明。

三、应用

在计算机体系中，对弗林分类法进行过说明，特别是其具体的内容都进行了阐述。而按弗林分类法中，SIMD用于对GPU体系的描述，但实际上，SIMD更倾向于数据的并行。但在CUDA范式中，其实是采用了NIVDA自己的SIMT结构，这个结构在前面也分析过了，这里不再赘述。在SIMT结构中，是由多个线程同时向不同的数据发出相同的操作指令，它更适合于NVIDA的GPU产品。
所以，目前广泛应用的AI体系的底层，特别是各种大模型的底层编程框架，仍然以CUDA为主，即以SIMT为主。就这一点，就保证了CUDA范式即异构编程的广泛性。
另外，在现在，云的迅猛发展以及各种设备的互联互通，使得异构编程成为了一种必然趋势。

四、例程

在AI和科学计算中，矩阵的运算是一个非常应用广泛的情况，下面看一个例子：

#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include <iostream>
#include <cassert>


__global__ void matrixAdd(float* A, float* B, float* C, int rows, int cols) {
	int row = blockIdx.y * blockDim.y + threadIdx.y;
	int col = blockIdx.x * blockDim.x + threadIdx.x;

	int id = row * cols + col;

	if (row < rows && col < cols) {
		C[id] = A[id] + B[id];
	}
}

int main() {
	int rows = 1 << 10; // 1024
	int cols = 1 << 10; // 1024
	int numElem = rows * cols;
	size_t size = numElem * sizeof(float);

	std::cout << "malloc : " << size / (1024 * 1024) << " MB" << std::endl;

	float *mA = new float[numElem];
	float *mB = new float[numElem];
	float *mC = new float[numElem];

	for (int i = 0; i < numElem; ++i) {
		mA[i] = rand() / (float)RAND_MAX;
		mB[i] = rand() / (float)RAND_MAX;
	}

	float *dmA, *dmB, *dmC;
	cudaMalloc((void**)&dmA, size);
	cudaMalloc((void**)&dmB, size);
	cudaMalloc((void**)&dmC, size);

	cudaMemcpy(dmA, mA, size, cudaMemcpyHostToDevice);
	cudaMemcpy(dmB, mB, size, cudaMemcpyHostToDevice);

	dim3 bSize(16, 16);
	dim3 gSize((cols + bSize.x - 1) / bSize.x,
		(rows + bSize.y - 1) / bSize.y);

	std::cout << "thread block num: [" << gSize.x << " * " << gSize.y
		<< "] , each block have[" << bSize.x << " * " << bSize.y
		<< "] threads " << std::endl;

	matrixAdd << <gSize, bSize >> >(dmA, dmB, dmC, rows, cols);

	cudaError_t err = cudaGetLastError();
	if (err != cudaSuccess) {
		std::cerr << "CUDA exec err: " << cudaGetErrorString(err) << std::endl;
		exit(EXIT_FAILURE);
	}

	cudaMemcpy(mC, dmC, size, cudaMemcpyDeviceToHost);

	cudaDeviceSynchronize();
	cudaDeviceProp prop;
	cudaGetDeviceProperties(&prop, 0);
	std::cout  << prop.name << "exec completed!" << std::endl;

	cudaFree(dmA);
	cudaFree(dmB);
	cudaFree(dmC);
	delete[] mA;
	delete[] mB;
	delete[] mC;

	cudaDeviceReset();

	std::cout << "The program has been executed successfully ！" << std::endl;
	return 0;
}