#include <iostream>
#include <cuda_runtime.h>
// 定义数组大小
const int arraySize = 1024;
// CUDA核函数,计算累加和
__global__ void sumArray(int* array, int* result) {
__shared__ int sharedData[arraySize]; // 共享内存,用于线程协作
int tid = threadIdx.x;
int blockOffset = blockIdx.x * blockDim.x;
// 将全局数组数据加载到共享内存中
sharedData[tid] = array[blockOffset + tid];
__syncthreads(); // 线程同步,等待所有线程加载完毕
// 进行累加
for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {
if (tid < stride) {
sharedData[tid] += sharedData[tid + stride];
}
__syncthreads(); // 线程同步,等待所有线程完成累加
}
// 累加结果写回全局内存
if (tid == 0) {
result[blockIdx.x] = sharedData[0];
}
}
int main() {
int hostArray[arraySize];
int* deviceArray;
int* deviceResult;
int result[arraySize];
// 初始化主机数组
for (int i = 0; i < arraySize; ++i) {
hostArray[i] = i;
}
// 在GPU上分配内存
cudaMalloc((void**)&deviceArray, sizeof(int) * arraySize);
cudaMalloc((void**)&deviceResult, sizeof(int) * arraySize);
// 将主机数据复制到GPU
cudaMemcpy(deviceArray, hostArray, sizeof(int) * arraySize, cudaMemcpyHostToDevice);
// 计算线程块和线程数
int blockSize = 256;
int numBlocks = (arraySize + blockSize - 1) / blockSize;
// 调用CUDA核函数
sumArray << <numBlocks, blockSize >> > (deviceArray, deviceResult);
// 将结果从GPU复制回主机
cudaMemcpy(result, deviceResult, sizeof(int) * numBlocks, cudaMemcpyDeviceToHost);
// 计算最终结果
int finalResult = 0;
for (int i = 0; i < numBlocks; ++i) {
finalResult += result[i];
}
std::cout << "累加和: " << finalResult << std::endl;
// 释放GPU内存
cudaFree(deviceArray);
cudaFree(deviceResult);
return 0;
}
cuda:使用共享内存
于 2023-09-11 16:35:01 首次发布