核函数 共享内存示例

### 核函数与共享内存使用 在CUDA编程中，核函数可以通过`__global__`关键字来定义，并能够在GPU上并行执行多个线程。当涉及到数据处理时，为了提高性能，通常会利用共享内存这一特性。共享内存在同一块SM（流多处理器）内的线程间提供快速访问速度。 对于如何在一个核函数内部有效地运用共享内存，下面给出一段示例代码，展示了两个浮点数数组相加的过程，在此过程中应用了共享内存优化[^4]： ```cpp #include <cuda.h> #include <stdio.h> #define N 1024 // 数组长度 #define THREADS_PER_BLOCK 256 // 定义全局核函数 __global__ void vecAddWithSharedMem(const float* A, const float* B, float* C){ extern __shared__ float sharedA[]; // 声明动态分配的共享内存 unsigned int tid = threadIdx.x; unsigned int i = blockIdx.x * blockDim.x + threadIdx.x; if(i<N){ // 将输入的数据加载到共享内存中 sharedA[tid] = A[i]; sharedA[blockDim.x + tid] = B[i]; // 确保所有线程都完成了对共享内存的写入操作后再读取 __syncthreads(); // 访问共享内存中的数据完成计算 C[i] = sharedA[tid] + sharedA[blockDim.x + tid]; } } int main(){ size_t bytes = N*sizeof(float); // 主机端指针声明 float h_A[N],h_B[N],h_C[N]; // 设备端指针声明及分配显存空间 float *d_A,*d_B,*d_C; cudaMalloc(&d_A,bytes); cudaMalloc(&d_B,bytes); cudaMalloc(&d_C,bytes); // 初始化主机端数据 for(int i=0;i<N;++i){ h_A[i]=rand()/(float)RAND_MAX; h_B[i]=rand()/(float)RAND_MAX; } // 数据传输至设备端 cudaMemcpy(d_A,h_A,bytes,cudaMemcpyHostToDevice); cudaMemcpy(d_B,h_B,bytes,cudaMemcpyHostToDevice); // 设置启动配置参数并调用核函数 dim3 block(THREADS_PER_BLOCK); dim3 grid((N+block.x-1)/block.x); vecAddWithSharedMem<<<grid,block,N*2*sizeof(float)>>>(d_A,d_B,d_C); // 获取结果回到主机端 cudaMemcpy(h_C,d_C,bytes,cudaMemcpyDeviceToHost); // 清理资源 cudaFree(d_A); cudaFree(d_B); cudaFree(d_C); return 0; } ``` 这段程序首先初始化了一个大小为N的一维向量，并将其复制到了GPU上的相应位置；接着设置了合适的网格尺寸和区块数量以匹配待解决问题规模；最后通过调用带有共享内存支持版本的矢量求和核函数实现了高效运算。