要求写一个reduceKernel 要求给出Kerne的完整调用:
1. 进行一维reduce 要求如下:
-
可以写一个最基础的,仅仅实现基础功能就行
-
使用share mem进行功能优化
-
使用shuffles指令完成block reduce操作
// 简单的实现,使用一个block完成Reduce操作 好蠢好蠢的代码
template<typename T>
__global__ void BlockReduce(const T* input, T* output, int num) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
for(int i = idx; i < num; i+= blockDim.x + gridDim.x) {
atomicAdd(output, input +i);
}
}
# 一个block完成reduce操作 使用sharemem blockSize = 256
__global__ void BlockReduceSharemem(const T* input, T* output, int num) {
int idx = threadIdx.x;
__share__ int sharemem[ThreadPerBlock];
sharemem[threadIdx.x] = 0;
__syncthreads();
for(int index = idx; index < num; index+=blockDim.x) {
sharemem[threadIdx.x] +=