cuda编程笔记（2.5）--简易的应用代码

矢量相加

#include <iostream>
#include <cuda_runtime.h>
#include "device_launch_parameters.h"__global__ void add(int* a, int* b, int* c,const int N) {int idx = blockIdx.x * blockDim.x + threadIdx.x;if(idx<N)c[idx] = a[idx] + b[idx];
}
const int N = 5;
int main() {int a[N] = { 1,2,3,4,5 }, b[N] = { 5,4,3,2,1 }, c[N] = { 0 };int* dev_a, * dev_b, * dev_c;cudaMalloc((void**)&dev_a, N * sizeof(int));cudaMalloc((void**)&dev_b, N * sizeof(int));cudaMalloc((void**)&dev_c, N * sizeof(int));cudaMemcpy(dev_a, a, N * sizeof(int), cudaMemcpyHostToDevice);cudaMemcpy(dev_b, b, N * sizeof(int), cudaMemcpyHostToDevice);cudaMemcpy(dev_c, c, N * sizeof(int), cudaMemcpyHostToDevice);add <<<N, 1 >>> (dev_a, dev_b, dev_c,N);cudaMemcpy(c, dev_c, N * sizeof(int), cudaMemcpyDeviceToHost);for (int i = 0; i < N; i++) {std::cout << c[i] << " ";}cudaDeviceSynchronize();return 0;
}

矢量点乘

#ifndef __CUDACC__
#define __CUDACC__
#endif
#include <iostream>
#include <cuda_runtime.h>
#include "device_launch_parameters.h"#define threadsPerBlock 256
const int Blocks = 32;
const int N = Blocks * threadsPerBlock;__global__ void dot(float* a, float* b, float* c) {__shared__ float cache[threadsPerBlock];int tid = threadIdx.x + blockIdx.x * blockDim.x;int cacheIndex = threadIdx.x;float temp = 0;if (tid < N) {temp = a[tid] * b[tid];}cache[cacheIndex] = temp;__syncthreads();// 并行归约求和for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {if (cacheIndex < stride) {cache[cacheIndex] += cache[cacheIndex + stride];}__syncthreads();}// 将每个 block 的结果写入全局内存if (cacheIndex == 0) {c[blockIdx.x] = cache[0];}
}int main() {// 1. 在主机分配内存float* a = new float[N];float* b = new float[N];// 初始化向量数据for (int i = 0; i < N; ++i) {a[i] = 1.0f;  // 举例，全1b[i] = 2.0f;  // 举例，全2}// 2. 设备内存指针float* dev_a, * dev_b, * dev_c;// 3. 分配设备内存cudaMalloc((void**)&dev_a, N * sizeof(float));cudaMalloc((void**)&dev_b, N * sizeof(float));cudaMalloc((void**)&dev_c, Blocks * sizeof(float));  // 每个 block 一个结果// 4. 复制数据到设备cudaMemcpy(dev_a, a, N * sizeof(float), cudaMemcpyHostToDevice);cudaMemcpy(dev_b, b, N * sizeof(float), cudaMemcpyHostToDevice);// 5. 启动核函数dot << <Blocks, threadsPerBlock >> > (dev_a, dev_b, dev_c);// 等待 GPU 完成cudaDeviceSynchronize();// 6. 从设备拷回部分和float* partial_sums = new float[Blocks];cudaMemcpy(partial_sums, dev_c, Blocks * sizeof(float), cudaMemcpyDeviceToHost);// 7. CPU 端归约float final_sum = 0;for (int i = 0; i < Blocks; ++i) {final_sum += partial_sums[i];}// 输出结果std::cout << "Dot product result: " << final_sum << std::endl;// 8. 释放内存delete[] a;delete[] b;delete[] partial_sums;cudaFree(dev_a);cudaFree(dev_b);cudaFree(dev_c);return 0;
}

__shared__ 是 CUDA 中的一个关键字，用于声明 线程块（block）内共享的内存变量，也就是所谓的 共享内存（shared memory）。

• 作用范围：只在同一个线程块（block）中可见。

• 作用对象：声明在 __device__ 或 __global__ 函数中的变量。

• 共享者：同一个 block 中的所有线程。

• 速度：非常快，接近于寄存器速度（比全局内存 __global__ 更快很多）。

• 作用域：每个 block 有自己独立的一份共享内存。

特性	描述
存储位置	GPU 的共享内存（每个 SM 独立的一块）
可见性	仅 block 内部所有线程可见
生命周期	线程块的生命周期（block 结束后释放）
分配方式	静态（指定大小）或动态（内核调用时分配）

限制	原因
不同 block 不能共享 __shared__ 变量	每个 block 拥有自己的共享内存
共享内存大小有限	每个 SM 共享内存一般是 48KB（可查 cudaDeviceProp.sharedMemPerBlock）
需手动同步	使用 __syncthreads() 保证所有线程写入后再读取

__syncthreads() 是 CUDA 中的一个 线程同步函数，用于在 同一个 block 中的所有线程之间进行同步。它确保：

所有线程都执行到这一点后，才会继续执行后面的代码。

 __syncthreads() 常用于：

• 共享内存读写同步，防止数据竞争。

• 实现线程间的数据依赖（例如：前缀和、归约等并行算法）。

• 保证共享内存中数据已经被所有线程写入或读取。

注意点	说明
只能在 __global__ 或 __device__ 函数中使用	因为它作用于 device 上的线程
只能在 block 内使用	不同 block 之间无法同步
所有线程都必须执行到这一点	否则会导致 死锁（hang）或错误
不支持条件内执行（有时）	若你在 if 内调用，要确保所有线程都会到达该分支或都不进入，否则可能死锁

 如果vs识别不出的话，加上这些宏定义即可

#ifndef __CUDACC__
#define __CUDACC__
#endif