我是 CUDA 和一般算法的新手。有人可以告诉我我这样做是否正确,或者是否有更好的方法。一个问题是代码的输入和输出应该在 GPU 上,这样主机和设备之间就不会发生内存复制。
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
#include <stdint.h>
#include <iostream>
#define TILE_WIDTH 8
__global__ void gpu_sumElements(int height, int width, float *in, float *out){
extern __shared__ float cache[];
int w = blockIdx.x * blockDim.x + threadIdx.x; // Col // width
int h = blockIdx.y * blockDim.y + threadIdx.y;
int index = h * width + w;
int cacheIndex = threadIdx.y * blockDim.x + threadIdx.x;
float temp = 0;
if ((w < width) && (h < height)){
temp += in[index];
//index += (height * width);
}
cache[cacheIndex] = temp;
__syncthreads();
int i = (blockDim.x * blockDim.y) / 2;
while (i != 0){
if (cacheIndex < i)
cache[cacheIndex] += cache[cacheIndex + i];
__syncthreads();
i /= 2;
}
if (cacheIndex == 0)
out[blockIdx.y * gridDim.x + blockIdx.x] = cache[0];
}
int main(){
// Initial Parameters
int width = 2363;
int height = 781;
float my_sum = 0;
int block_height = (height - 1) / TILE_WIDTH + 1;
int block_width = (width - 1) / TILE_WIDTH + 1;
dim3 dimGrid(block_width, block_height, 1);
dim3 dimBlock(TILE_WIDTH, TILE_WIDTH, 1);
int sharedMemSize = TILE_WIDTH * TILE_WIDTH * sizeof(float);
// Initialize host arrays
float *test_array = new float[height * width];
float *out = new float[height * width];
for (int i = 0; i < (height * width); i++)
test_array[i] = 1.0f;
// Initialize device arrays
float *gpu_temp_array;
float *gpu_out;
cudaMalloc((void **)&gpu_temp_array, (height * width * sizeof(float)));
cudaMalloc((void **)&gpu_out, (height * width * sizeof(float)));
cudaMemcpy(gpu_out, test_array, (height * width * sizeof(float)), cudaMemcpyHostToDevice);
// Copy these, need them elsewhere
float sum_height = height;
float sum_width = width ;
dim3 sum_dimGrid = dimGrid;
int i = (height * width);
// Launch kernel, get & print results
while (i != 0){
gpu_sumElements<<<sum_dimGrid, dimBlock, sharedMemSize>>>(sum_height, sum_width, gpu_out, gpu_temp_array);
cudaMemcpy(gpu_out, gpu_temp_array, (sum_height * sum_width * sizeof(float)), cudaMemcpyDeviceToDevice);
cudaMemset(gpu_temp_array, 0, (height * width * sizeof(float)));
sum_height = ceil(sum_height/TILE_WIDTH);
sum_width = ceil(sum_width/TILE_WIDTH);;
sum_dimGrid.x = (sum_width - 1) / TILE_WIDTH + 1;
sum_dimGrid.y = (sum_height - 1) / TILE_WIDTH + 1;
i /= TILE_WIDTH*TILE_WIDTH;
}
cudaMemcpy(out, gpu_out, (height * width * sizeof(float)), cudaMemcpyDeviceToHost);
std::cout << out[0] << std::endl << std::endl;
delete[] test_array;
delete[] out;
cudaFree(gpu_out);
cudaFree(gpu_temp_array);
system("pause");
return 0;
}
最佳答案
一般来说,使用多个内核启动来并行缩减来生成一个(最终)结果通常是没有必要的。 cuda sample code 详细记录了生成组织良好的并行缩减的过程,该过程仅需要针对任意数据大小启动两次内核。和 accompanying PDF .
要创建仅使用单个内核启动的并行缩减,至少有两种常见方法:
使用所谓的“减少线程围栏”方法。这也在 CUDA sample code 中得到了体现。 。在这种方法中,最终的缩减阶段是通过跟踪“内核排出”来执行的。具体来说,每个线程 block 在完成其工作时更新“完成计数”变量(原子地)。由于启动的线程 block 的数量是已知的,因此线程 block 可以确定它是否是最后一个完成的线程 block 。如果是,则该线程 block 将其他线程 block 产生的所有中间结果相加,这些结果现在被写入全局内存。 “threadfence”绰号是因为每个线程 block 必须确保其部分结果在退出之前在全局内存中可用(使用 threadfence intrinsic )。该方法可以处理“任意”减少。
每个线程 block (其中有一个线程)atomically使用其自己的部分结果更新最终的内核范围结果。这仅对于提供相应原子函数的约简而言才可以方便地实现,例如求和、求最大值、求最小值等
上述任何一种方法都将受益于 CUDA 并行缩减示例代码中涵盖的基本技术,特别是将线程 block 的数量减少到仍允许充分利用 GPU 的最小值。这种优化允许最小数量的原子操作。考虑到这些优化,与相应的 2 内核或多内核减少相比,减少可以更快,并且“更简单”(例如,单个内核调用,无需对中间结果进行太多主机管理)。
关于c++ - CUDA:使用线性化 2D 共享内存的数组中所有元素的总和,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/25415239/