规约算法原理可参照上一篇关于规约算法求解数组和的博客,此处不再详细介绍,直接给出代码实现过程及注释。
#include <stdio.h>
#include <stdlib.h>
#define N 1000 // 数组大小
__global__ void findMax(int *array, int *maxValue, int *maxIndex) {
__shared__ int s_maxValue;//保证每一个块在实现执行此函数时都会首先申请一块共享内存用来存储最大值
__shared__ int s_maxIndex;//保证每一个块在实现执行此函数时都会首先申请一块共享内存用来存储最大值索引
int tid = threadIdx.x;
int idx = blockIdx.x * blockDim.x + threadIdx.x;
// 初始化共享内存
if (tid == 0) {
s_maxValue = array[0];
s_maxIndex = 0;
}
__syncthreads();
// 在每个线程块内找到局部最大值及其索引(注意是线程块内部)
for (int i = idx; i < N; i += blockDim.x * gridDim.x) {
if (array[i] > s_maxValue) {
s_maxValue = array[i];
s_maxIndex = i;
}
}
__syncthreads();
// 由于块间共享内存的数据无法实现共享,因此此处需要将每个线程块的最大值和索引都写回全局内存
if (tid == 0) {
maxValue[blockIdx.x] = s_maxValue;
maxIndex[blockIdx.x] = s_maxIndex;
}
}
int main() {
int array[N]; // 定义数组
int *dev_array, *dev_maxValue, *dev_maxIndex; // 定义CUDA变量
int maxValue[N / 256], maxIndex[N / 256]; // 存储每个线程块的最大值和索引
// 在设备上分配内存
cudaMalloc((void**)&dev_array, N * sizeof(int));
cudaMalloc((void**)&dev_maxValue, (N / 256) * sizeof(int));
cudaMalloc((void**)&dev_maxIndex, (N / 256) * sizeof(int));
// 初始化数组
for (int i = 0; i < N; i++) {
array[i] = rand() % 1000; // 使用更大的范围生成随机数
}
// 将数组拷贝到设备
cudaMemcpy(dev_array, array, N * sizeof(int), cudaMemcpyHostToDevice);
// 调用内核函数
findMax<<<N / 256, 256>>>(dev_array, dev_maxValue, dev_maxIndex);
// 将结果从设备拷贝回主机
cudaMemcpy(maxValue, dev_maxValue, (N / 256) * sizeof(int), cudaMemcpyDeviceToHost);
cudaMemcpy(maxIndex, dev_maxIndex, (N / 256) * sizeof(int), cudaMemcpyDeviceToHost);
// 在主机端找到全局最大值及其对应位置
int globalMaxValue = -1, globalMaxIndex = -1;
for (int i = 0; i < N / 256; i++) {
if (maxValue[i] > globalMaxValue) {
globalMaxValue = maxValue[i];
globalMaxIndex = maxIndex[i];
}
}
printf("数组最大值为: %d, 位于索引位置: %d\n", globalMaxValue, globalMaxIndex);
// 释放设备上分配的内存
cudaFree(dev_array);
cudaFree(dev_maxValue);
cudaFree(dev_maxIndex);
return 0;
}