当前位置：首页 > wzjs >正文

湖州网站建设湖州手机购物平台

wzjs 2025/9/3 22:05:27

湖州网站建设湖州,手机购物平台,天空影院手机免费观看在线,免费网站建设网站推广在第3章讨论过的并行归约问题 5.3.1 使用共享内存进行并行归约 reduceGmem – 使用全局内存作为基准 reduceSmem – 使用共享内存 #include <cuda_runtime.h> #include <stdio.h> #include "../common/common.h" #include <iostream>#define DIM…

在第3章讨论过的并行归约问题

5.3.1 使用共享内存进行并行归约
reduceGmem – 使用全局内存作为基准
reduceSmem – 使用共享内存

#include <cuda_runtime.h>
#include <stdio.h>
#include "../common/common.h"
#include <iostream>#define DIM 128int recursiveReduce(int *data, int const size){if (size == 1) return data[0];int const stride = size /2;for (int i = 0; i < stride; i ++){data[i] += data[i + stride];}return recursiveReduce( data, stride);
}__global__ void warmup( int *g_idata, int *g_odata, unsigned int n){unsigned int tid  = threadIdx.x;int *idata = g_idata + blockIdx.x * blockDim.x;unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;if (idx >= n) return;if (blockDim.x >= 1024 &&  tid < 512) idata[tid] += idata[tid+ 512];__syncthreads();if (blockDim.x >= 512 &&  tid < 256) idata[tid] += idata[tid+ 256];__syncthreads();if (blockDim.x >= 256 &&  tid < 128) idata[tid] += idata[tid+ 128];__syncthreads();if (blockDim.x >= 128 &&  tid < 64) idata[tid] += idata[tid+ 64];__syncthreads();if (tid < 32){volatile int *vmem  = idata;vmem[tid] += vmem[tid + 32];vmem[tid] += vmem[tid + 16];vmem[tid] += vmem[tid +  8];vmem[tid] += vmem[tid +  4];vmem[tid] += vmem[tid +  2];vmem[tid] += vmem[tid +  1];}if  (tid == 0){ g_odata[blockIdx.x] = idata[0];}
}__global__ void reduceGmem( int *g_idata, int *g_odata, unsigned int n){unsigned int tid  = threadIdx.x;int *idata = g_idata + blockIdx.x * blockDim.x;unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;if (idx >= n) return;if (blockDim.x >= 1024 &&  tid < 512) idata[tid] += idata[tid+ 512];__syncthreads();if (blockDim.x >= 512 &&  tid < 256) idata[tid] += idata[tid+ 256];__syncthreads();if (blockDim.x >= 256 &&  tid < 128) idata[tid] += idata[tid+ 128];__syncthreads();if (blockDim.x >= 128 &&  tid < 64) idata[tid] += idata[tid+ 64];__syncthreads();if (tid < 32){volatile int *vmem  = idata;vmem[tid] += vmem[tid + 32];vmem[tid] += vmem[tid + 16];vmem[tid] += vmem[tid +  8];vmem[tid] += vmem[tid +  4];vmem[tid] += vmem[tid +  2];vmem[tid] += vmem[tid +  1];}if  (tid == 0){ g_odata[blockIdx.x] = idata[0];}
}__global__ void reduceSmem(int *g_idata, int *g_odata, unsigned int n){__shared__ int smem[DIM];unsigned int tid  = threadIdx.x;// convert global data pointer to local pointerint *idata = g_idata + blockIdx.x * blockDim.x;unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;if (idx >= n) return;//set to smem by each threadssmem[tid] = idata[tid];__syncthreads();if (blockDim.x >= 1024 &&  tid < 512) smem[tid] += smem[tid+ 512];__syncthreads();if (blockDim.x >= 512 &&  tid < 256) smem[tid] += smem[tid+ 256];__syncthreads();if (blockDim.x >= 256 &&  tid < 128) smem[tid] += smem[tid+ 128];__syncthreads();if (blockDim.x >= 128 &&  tid < 64) smem[tid] += smem[tid+ 64];__syncthreads();if (tid < 32){volatile int *vsmem  = smem;vsmem[tid] += vsmem[tid + 32];vsmem[tid] += vsmem[tid + 16];vsmem[tid] += vsmem[tid +  8];vsmem[tid] += vsmem[tid +  4];vsmem[tid] += vsmem[tid +  2];vsmem[tid] += vsmem[tid +  1];}if  (tid == 0){ g_odata[blockIdx.x] = smem[0];}
}int main(int argc , char **argv)
{printf("%s starting\n", argv[0]);int dev = 0;cudaDeviceProp deviceprop;CHECK(cudaGetDeviceProperties(&deviceprop,dev));printf("Using Device %d : %s\n", dev, deviceprop.name);int size = 1 << 24;int blocksize = 512;if (argc > 1){blocksize = atoi(argv[1]);}dim3 block(DIM, 1);  // 1ddim3 grid ((size + block.x - 1) / block.x, 1);size_t nBytes = size  * sizeof(int);int * h_idata = (int*) malloc(nBytes);int * h_odata = (int*) malloc( grid.x * sizeof(int));  //you duoshao ge blockint * temp = (int*) malloc(nBytes);//initial the arrayfor (int i = 0 ; i < size;i++){h_idata[i] = (int)(rand() & 0xff);}int sum = 0;for (int i = 0 ; i < size;i++){sum += h_idata[i];}printf("sum value is : %d\n", sum);memcpy(temp, h_idata, nBytes);int gpu_sum = 0;int *d_idata = NULL;int *d_odata = NULL;cudaMalloc((void**)&d_idata, nBytes);cudaMalloc((void**)&d_odata, grid.x * sizeof(int));//cpu sumTimer timer;timer.start();int cpu_sum = recursiveReduce(temp, size);timer.stop();float elapsedTime = timer.elapsedms();printf("cpu reduce time: %f,  sum: %d\n", elapsedTime, cpu_sum);//gpu sumcudaMemcpy(d_idata, h_idata, nBytes, cudaMemcpyHostToDevice);cudaDeviceSynchronize();timer.start();warmup<<<grid.x, block>>>(d_idata, d_odata, size);cudaDeviceSynchronize(); timer.stop();float elapsedTime1 = timer.elapsedms();cudaMemcpy(h_odata, d_odata, grid.x * sizeof(int),cudaMemcpyDeviceToHost);gpu_sum = 0;for (int i = 0; i < grid.x; i ++){gpu_sum += h_odata[i];}printf("warm up reduce time: %f,  sum: %d\n", elapsedTime1, gpu_sum);//gpu sumcudaMemcpy(d_idata, h_idata, nBytes, cudaMemcpyHostToDevice);cudaDeviceSynchronize();timer.start();reduceGmem<<<grid.x, block>>>(d_idata, d_odata, size);cudaDeviceSynchronize(); timer.stop();elapsedTime1 = timer.elapsedms();cudaMemcpy(h_odata, d_odata, grid.x * sizeof(int),cudaMemcpyDeviceToHost);gpu_sum = 0;for (int i = 0; i < grid.x ; i ++){gpu_sum += h_odata[i];}printf("reduceGmem gpu reduce time: %f,  sum: %d, gird ,block (%d %d)\n", elapsedTime1, gpu_sum, grid.x , block.x);//gpu sumcudaMemcpy(d_idata, h_idata, nBytes, cudaMemcpyHostToDevice);cudaDeviceSynchronize();timer.start();reduceSmem<<<grid.x, block>>>(d_idata, d_odata, size);cudaDeviceSynchronize(); timer.stop();elapsedTime1 = timer.elapsedms();cudaMemcpy(h_odata, d_odata, grid.x * sizeof(int),cudaMemcpyDeviceToHost);gpu_sum = 0;for (int i = 0; i < grid.x ; i ++){gpu_sum += h_odata[i];}printf("reduceSmem gpu reduce time: %f,  sum: %d, gird ,block (%d %d)\n", elapsedTime1, gpu_sum, grid.x , block.x);cudaFree(d_idata);cudaFree(d_odata);cudaDeviceReset();free(h_idata);free(h_odata);free(temp);return 0;
}

通过nsys profile 程序：
nsys profile --stats=true reduce.exe
输出：

Time (%)  Total Time (ns)  Instances  Avg (ns)  Med (ns)  Min (ns)  Max (ns)  StdDev (ns)                   Name--------  ---------------  ---------  --------  --------  --------  --------  -----------  --------------------------------------39.8           134721          1  134721.0  134721.0    134721    134721          0.0  warmup(int *, int *, unsigned int)38.0           128385          1  128385.0  128385.0    128385    128385          0.0  reduceGmem(int *, int *, unsigned int)22.2            75040          1   75040.0   75040.0     75040     75040          0.0  reduceSmem(int *, int *, unsigned int)

5.3.2 展开

展开的核函数以及调用：


__global__ void reduceSmemUnroll(int *g_idata, int *g_odata, unsigned int n){__shared__ int smem[DIM];unsigned int tid  = threadIdx.x;// convert global data pointer to local pointerint *idata = g_idata + blockIdx.x * blockDim.x;unsigned int idx = blockIdx.x * blockDim.x * 4 + threadIdx.x;//unrolling 4 blocksint tmpSum = 0;if (idx + 3 * blockDim.x <= n){int a1 = g_idata[idx];int a2 = g_idata[idx + blockDim.x];int a3 = g_idata[idx + 2 * blockDim.x];int a4 = g_idata[idx + 3 * blockDim.x];tmpSum = a1 + a2 + a3 + a4;}//set to smem by each threadssmem[tid] = tmpSum;__syncthreads();if (blockDim.x >= 1024 &&  tid < 512) smem[tid] += smem[tid+ 512];__syncthreads();if (blockDim.x >= 512 &&  tid < 256) smem[tid] += smem[tid+ 256];__syncthreads();if (blockDim.x >= 256 &&  tid < 128) smem[tid] += smem[tid+ 128];__syncthreads();if (blockDim.x >= 128 &&  tid < 64) smem[tid] += smem[tid+ 64];__syncthreads();if (tid < 32){volatile int *vsmem  = smem;vsmem[tid] += vsmem[tid + 32];vsmem[tid] += vsmem[tid + 16];vsmem[tid] += vsmem[tid +  8];vsmem[tid] += vsmem[tid +  4];vsmem[tid] += vsmem[tid +  2];vsmem[tid] += vsmem[tid +  1];}if  (tid == 0){ g_odata[blockIdx.x] = smem[0];}
}// 调用
cudaMemcpy(d_idata, h_idata, nBytes, cudaMemcpyHostToDevice);cudaDeviceSynchronize();timer.start();reduceSmemUnroll<<<grid.x /4 , block>>>(d_idata, d_odata, size);cudaDeviceSynchronize(); timer.stop();elapsedTime1 = timer.elapsedms();cudaMemcpy(h_odata, d_odata, grid.x /4 * sizeof(int),cudaMemcpyDeviceToHost);gpu_sum = 0;for (int i = 0; i < grid.x / 4 ; i ++){gpu_sum += h_odata[i];}printf("reduceSmemUnroll gpu reduce time: %f,  sum: %d, gird ,block (%d %d)\n", elapsedTime1, gpu_sum, grid.x / 4, block.x);

nsys输出：

Time (%)  Total Time (ns)  Instances  Avg (ns)  Med (ns)  Min (ns)  Max (ns)  StdDev (ns)                      Name--------  ---------------  ---------  --------  --------  --------  --------  -----------  --------------------------------------------34.6           135329          1  135329.0  135329.0    135329    135329          0.0  warmup(int *, int *, unsigned int)33.0           129056          1  129056.0  129056.0    129056    129056          0.0  reduceGmem(int *, int *, unsigned int)      25.6            99968          1   99968.0   99968.0     99968     99968          0.0  reduceSmem(int *, int *, unsigned int)6.7            26335          1   26335.0   26335.0     26335     26335          0.0  reduceSmemUnroll(int *, int *, unsigned int)

5.3.3 动态共享内存

//动态声明
extern __shared__ int smem[];//调用
reduceSmemUnrollDyn<<<grid.x /4 , block, DIM * sizeof(int)>>>(d_idata, d_odata, size);

发现用动态分配共享内存实现的核函数和用静态分配共享内存实现的核函数之间没有显著的差异。

Time (%)  Total Time (ns)  Instances  Avg (ns)  Med (ns)  Min (ns)  Max (ns)  StdDev (ns)                       Name--------  ---------------  ---------  --------  --------  --------  --------  -----------  -----------------------------------------------35.2           135009          1  135009.0  135009.0    135009    135009          0.0  warmup(int *, int *, unsigned int)33.5           128576          1  128576.0  128576.0    128576    128576          0.0  reduceGmem(int *, int *, unsigned int)19.5            74753          1   74753.0   74753.0     74753     74753          0.0  reduceSmem(int *, int *, unsigned int)5.9            22752          1   22752.0   22752.0     22752     22752          0.0  reduceSmemUnroll(int *, int *, unsigned int)5.9            22752          1   22752.0   22752.0     22752     22752          0.0  reduceSmemUnrollDyn(int *, int *, unsigned int)