在中介绍了使用一个包含N个线程的线程块和共享内存进行数组归约求和,
基本思路:
定义M个包含N个线程的线程块时(NThreadX = ((NX + ThreadX - 1) / ThreadX)),全局线程索引需使用tid = blockIdx.x * blockDim.x + threadIdx.x,而在每个线程块中局部线程索引是i = threadIdx.x,
每个线程块只计算一部分求和,求和结果保存在该线程块中的共享内存数组0号元素中,线程结束后将该值赋给对应全局数组(blockIdx.x * blockDim.x)元素中,最后在CPU端使用循环将每个线程块所求和相加,即得到最后结果。
代码如下:
#pragma once#include "cuda_runtime.h"#include "device_launch_parameters.h"#include "device_functions.h"#includeusing namespace std;const int NX = 10240; //数组长度const int ThreadX = 256; //线程块大小//使用shared memory和多个线程块__global__ void d_SharedMemoryTest(double *para){ int i = threadIdx.x; //该线程块中线程索引 int tid = blockIdx.x * blockDim.x + threadIdx.x; //M个包含N个线程的线程块中相对应全局内存数组的索引(全局线程) __shared__ double s_Para[ThreadX]; //定义固定长度(线程块长度)的共享内存数组 if (tid < NX) //判断全局线程小于整个数组长度NX,防止数组越界 s_Para[i] = para[tid]; //将对应全局内存数组中一段元素的值赋给共享内存数组 __syncthreads(); //(红色下波浪线提示由于VS不识别,不影响运行)同步,等待所有线程把自己负责的元素载入到共享内存再执行下面代码 for (int index = 1; index < blockDim.x; index *= 2) //归约求和 { __syncthreads(); if (i % (2 * index) == 0) { s_Para[i] += s_Para[i + index]; } } if (i == 0) //求和完成,总和保存在共享内存数组的0号元素中 para[blockIdx.x * blockDim.x + i] = s_Para[i]; //在每个线程块中,将共享内存数组的0号元素赋给全局内存数组的对应元素,即线程块索引*线程块维度+i(blockIdx.x * blockDim.x + i)}//使用shared memory和多个线程块void s_ParallelTest(){ double *Para; cudaMallocManaged((void **)&Para, sizeof(double) * NX); //统一内存寻址,CPU和GPU都可以使用 double ParaSum = 0; for (int i = 0; i > > (Para); //调用核函数(M个包含N个线程的线程块) cudaDeviceSynchronize(); //同步 for (int i=0; i
结果如下(CPU和GPU结果一致):