我的电脑没有装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和B
for(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*2
return 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):
// 初始化点积结果为0
dotProductResult := 0
// 遍历两个向量的每个元素并相乘累加
for i from 0 to N-1 do
dotProductResult := 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; // 假设数组长度为1024
const 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)指定每个块需要的共享内存大小，这里每个线程计算一个部分和，然后存入共享内存，所以我们需要为每个块分配足够大的共享内存来存储这些部分和。
？
