本文主要是介绍CUDA C编程:第一个程序 向量相加点积,希望对大家解决编程问题提供一定的参考价值,需要的开发者们随着小编来一起学习吧!
我的电脑没有装CUDA,所以使用租了带GPU的云服务器,然后使用vscode SSH远程连接云服务器。云GPU使用的是智星云,0.8元/h。
智星云
可以使用nvcc --version查看系统中安装的CUDA版本。
然后写第一个CUDA程序,两个向量相加结果给到第三个向量
#include <cuda_runtime.h>
#include <iostream>#define CHECK(call) \
{ \const cudaError_t error = call; \if (error != cudaSuccess) { \std::cerr << "Error: " << __FILE__ << ", line " << __LINE__ << ": " \<< cudaGetErrorString(error) << std::endl; \exit(1); \} \
}__global__ void addArrays(const int *A, const int *B, int *C, int N) {int idx = blockIdx.x * blockDim.x + threadIdx.x;if (idx < N)C[idx] = A[idx] + B[idx];
}int main() {const int N = 100; // 数组大小int A[N], B[N], C[N];// 初始化数组A和Bfor(int i = 0; i < N; ++i) {A[i] = i;B[i] = i * 2;}int *d_A, *d_B, *d_C;// 分配GPU内存CHECK(cudaMalloc((void**)&d_A, N * sizeof(int)));CHECK(cudaMalloc((void**)&d_B, N * sizeof(int)));CHECK(cudaMalloc((void**)&d_C, N * sizeof(int)));// 将数据从主机复制到设备CHECK(cudaMemcpy(d_A, A, N * sizeof(int), cudaMemcpyHostToDevice));CHECK(cudaMemcpy(d_B, B, N * sizeof(int), cudaMemcpyHostToDevice));// 调用核函数addArrays<<<10, 10>>>(d_A, d_B, d_C, N);// 同步以确保核函数执行完成cudaDeviceSynchronize();// 将结果从设备复制回主机CHECK(cudaMemcpy(C, d_C, N * sizeof(int), cudaMemcpyDeviceToHost));// 释放GPU内存CHECK(cudaFree(d_A));CHECK(cudaFree(d_B));CHECK(cudaFree(d_C));// 输出结果for(int i = 0; i < N; ++i)std::cout << C[i] << " "; // 应该输出 i + i*2return 0;
}
nvcc -o add add.cu编译程序
./add运行程序
程序说明
#include <cuda_runtime.h>
引入cuda运行时环境
#define CHECK(call) \
{ \const cudaError_t error = call; \if (error != cudaSuccess) { \std::cerr << "Error: " << __FILE__ << ", line " << __LINE__ << ": " \<< cudaGetErrorString(error) << std::endl; \exit(1); \} \
}
用来提供CUDA报错信息的宏,用CHECK宏嵌套每一个将要调用的函数,便于调试。
#define CHECK(call) 定义了一个名为CHECK的宏,它接受一个参数call,这个参数是想检查的CUDA API调用。接下来的花括号 { ... } 包围了宏展开后将要执行的代码块。
const cudaError_t error=call;执行传入的CUDA API调用(即call),并将其返回的错误状态保存在变量error中。
if(error!=cudaSuccess){...}:检查error是否等于cudaSuccess,这是CUDA中表示操作成功的常量。如果不等于(即操作失败),则执行大括号内的错误处理代码。
std::cerr<< "Error: " <<__FILE__<<", line "<<__LINE__<<": "<<cudaGetErrorString(error)<< std::endl; 这行代码打印错误信息到标准错误输出。包括了出错的文件名(由__FILE__宏提供)、行号(由__LINE__宏提供),以及通过cudaGetErrorString(error)获取的错误描述字符串。exit(1); 如果确实发生了错误,程序会调用exit(1)立即终止,返回码1通常表示异常终止。
__global__ void addArrays(const int *A, const int *B, int *C, int N) {int idx = blockIdx.x * blockDim.x + threadIdx.x;if (idx < N)C[idx] = A[idx] + B[idx];
}
__global__ 是一个关键字,用于声明一个在GPU上执行的函数,也称为全局函数或内核函数。这些函数由主机(CPU)调用,但在设备(GPU)上的多个线程并行执行。
void addArrays(const int *A,const int *B,int *C, int N)定义了内核函数addArrays。
const int *A 和 const int *B指向输入数组A和B的指针,在内核中只读。
int *C输出数组C的指针,存放A和B对应元素的和。
N:需要相加的元素个数。
int idx = blockIdx.x * blockDim.x + threadIdx.x; 计算当前线程的全局索引 idx。
这里是CUDA线程组织方式的一个体现:
blockIdx.x 是当前线程所在的块(block)在网格(grid)中的x轴索引。
blockDim.x 是每个块中线程的数量(块的尺寸)在x轴方向。
threadIdx.x 是当前线程在块内的x轴索引。 通过这样的计算,每个线程都能知道自己在整个计算任务中的唯一位置,从而决定应该处理哪个数组元素。
if (idx < N) 是一个边界检查,确保线程不会访问超过数组界限。因为CUDA会为整个网格启动比实际需要更多的线程以充分利用硬件资源,所以这种检查是必要的。
C[idx] = A[idx] + B[idx]; 如果索引idx在有效范围内,这个语句就执行数组A和B中相应位置的元素相加,并将结果存储到数组C的相同位置。
(这个地方还是没怎么看懂)。
cudaMalloc((void**)&d_A, N * sizeof(int))
给设备分配N个int类型的内存,使用指针变量d_A指示。
cudaMemcpy(d_A, A, N * sizeof(int), cudaMemcpyHostToDevice)
内存拷贝,从Host拷贝到Device。A数组赋值给d_A数组。
计算两个数组的点积
理解CUDA内核调用中的<<< >>>语法。
向量点积计算伪代码
function dotProduct(A, B, N):// 初始化点积结果为0dotProductResult := 0// 遍历两个向量的每个元素并相乘累加for i from 0 to N-1 dodotProductResult := dotProductResult + (A[i] * B[i])
假设有一个简单的CUDA内核函数,用于计算两个数组的点积,并将结果存储在一个变量中。
__global__ void dotProductKernel(const float* A, const float* B,
float* result, int N) {extern __shared__ float partialSums[];unsigned int tid = threadIdx.x;unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;float sum = 0.0f;if (i < N) {for (unsigned int j = 0; j < N; j++) {sum += A[j * blockDim.x + tid] * B[j * blockDim.x + tid];}}partialSums[tid] = sum;__syncthreads(); // 确保所有线程完成上面的计算// 如果是块内的第一个线程,则累加所有部分和if (tid == 0) {for (unsigned int i = 0; i < blockDim.x; i++) {*result += partialSums[i];}}
}
在这个内核中,我们想要计算两个长度为N的一维数组A和B的点积。为了简化说明,我们忽略了一些优化(如减少共享内存的银行冲突),专注于展示如何调用这个内核。
现在,让我们看看如何使用<<< >>>来调用这个内核函数,并配置执行环境:
int main() {const int N = 1024; // 假设数组长度为1024const int blockSize = 256; // 每个块包含256个线程const int gridSize = (N + blockSize - 1) / blockSize; // 计算所需块的数量float *d_A, *d_B, *d_result;float h_result = 0.0f;float h_A[N], h_B[N]; // 主机端的数组// 初始化h_A和h_B数组,略...// 分配和复制数据到GPU,略...// 调用内核函数,注意 <<<gridSize, blockSize, sizeof(float)*blockSize>>> 的用法dotProductKernel<<<gridSize, blockSize, blockSize * sizeof(float)>>>(d_A, d_B, &h_result, N);// 将结果从GPU复制回CPU,略...// 释放GPU资源,略...return 0;
}
dotProductKernel<<<gridSize, blockSize, blockSize * sizeof(float)>>>(d_A, d_B, &h_result, N);是内核调用的实例。
gridSize和blockSize分别定义了执行该内核的网格和块的大小。
gridSize = 4因为1024个元素,每个块处理256个,共需要4个块和blockSize = 256。
第三个参数blockSize * sizeof(float)指定每个块需要的共享内存大小,这里每个线程计算一个部分和,然后存入共享内存,所以我们需要为每个块分配足够大的共享内存来存储这些部分和。
?
这篇关于CUDA C编程:第一个程序 向量相加点积的文章就介绍到这儿,希望我们推荐的文章对编程师们有所帮助!