我的电脑没有装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)指定每个块需要的共享内存大小，这里每个线程计算一个部分和，然后存入共享内存，所以我们需要为每个块分配足够大的共享内存来存储这些部分和。

？