include_starpu/skepu/src/reduce_kernels.h

Go to the documentation of this file.

#ifndef REDUCE_KERNELS_H
#define REDUCE_KERNELS_H

#ifdef SKEPU_OPENCL

#include <string>

namespace skepu
{

static std::string ReduceKernel_CL(
"__kernel void ReduceKernel_KERNELNAME(__global TYPE* input, __global TYPE* output, unsigned int n, __local TYPE* sdata)\n"
"{\n"
"    unsigned int blockSize = get_local_size(0);\n"
"    unsigned int tid = get_local_id(0);\n"
"    unsigned int i = get_group_id(0)*blockSize + get_local_id(0);\n"
"    unsigned int gridSize = blockSize*get_num_groups(0);\n"
"    TYPE result = 0;\n"
"    if(i < n)\n"
"    {\n"
"        result = input[i];\n"
"        i += gridSize;\n"
"    }\n"
"    while(i < n)\n"
"    {\n"
"        result = FUNCTIONNAME(result, input[i], (TYPE)0);\n"
"        i += gridSize;\n"
"    }\n"
"    sdata[tid] = result;\n"
"    barrier(CLK_LOCAL_MEM_FENCE);\n"
"    if(blockSize >= 512) { if (tid < 256 && tid + 256 < n) { sdata[tid] = FUNCTIONNAME(sdata[tid], sdata[tid + 256], (TYPE)0); } barrier(CLK_LOCAL_MEM_FENCE); }\n"
"    if(blockSize >= 256) { if (tid < 128 && tid + 128 < n) { sdata[tid] = FUNCTIONNAME(sdata[tid], sdata[tid + 128], (TYPE)0); } barrier(CLK_LOCAL_MEM_FENCE); }\n"
"    if(blockSize >= 128) { if (tid <  64 && tid +  64 < n) { sdata[tid] = FUNCTIONNAME(sdata[tid], sdata[tid +  64], (TYPE)0); } barrier(CLK_LOCAL_MEM_FENCE); }\n"
"    if(blockSize >=  64) { if (tid <  32 && tid +  32 < n) { sdata[tid] = FUNCTIONNAME(sdata[tid], sdata[tid +  32], (TYPE)0); } barrier(CLK_LOCAL_MEM_FENCE); }\n"
"    if(blockSize >=  32) { if (tid <  16 && tid +  16 < n) { sdata[tid] = FUNCTIONNAME(sdata[tid], sdata[tid +  16], (TYPE)0); } barrier(CLK_LOCAL_MEM_FENCE); }\n"
"    if(blockSize >=  16) { if (tid <   8 && tid +   8 < n) { sdata[tid] = FUNCTIONNAME(sdata[tid], sdata[tid +   8], (TYPE)0); } barrier(CLK_LOCAL_MEM_FENCE); }\n"
"    if(blockSize >=   8) { if (tid <   4 && tid +   4 < n) { sdata[tid] = FUNCTIONNAME(sdata[tid], sdata[tid +   4], (TYPE)0); } barrier(CLK_LOCAL_MEM_FENCE); }\n"
"    if(blockSize >=   4) { if (tid <   2 && tid +   2 < n) { sdata[tid] = FUNCTIONNAME(sdata[tid], sdata[tid +   2], (TYPE)0); } barrier(CLK_LOCAL_MEM_FENCE); }\n"
"    if(blockSize >=   2) { if (tid <   1 && tid +   1 < n) { sdata[tid] = FUNCTIONNAME(sdata[tid], sdata[tid +   1], (TYPE)0); } barrier(CLK_LOCAL_MEM_FENCE); }\n"
"    if(tid == 0)\n"
"    {\n"
"        output[get_group_id(0)] = sdata[tid];\n"
"    }\n"
"}\n"
);

}

#endif

#ifdef SKEPU_CUDA

namespace skepu
{

template<typename T, typename BinaryFunc>
__global__ void ReduceKernel_CU(BinaryFunc reduceFunc, T* input, T* output, unsigned int n)
{
    //A bit ugly
    extern __shared__ char _sdata[];
    T* sdata = reinterpret_cast<T*>(_sdata);

    unsigned int blockSize = blockDim.x;
    unsigned int tid = threadIdx.x;
    unsigned int i = blockIdx.x * blockSize + tid;
    unsigned int gridSize = blockSize*gridDim.x;
    T result = 0;

    if(i < n)
    {
        result = input[i];
        i += gridSize;
    }

    while(i < n)
    {
        result = reduceFunc.CU(result, input[i]);
        i += gridSize;
    }

    sdata[tid] = result;

    __syncthreads();

    if(blockSize >= 512) { if (tid < 256 && tid + 256 < n) { sdata[tid] = reduceFunc.CU(sdata[tid], sdata[tid + 256]); } __syncthreads(); }
    if(blockSize >= 256) { if (tid < 128 && tid + 128 < n) { sdata[tid] = reduceFunc.CU(sdata[tid], sdata[tid + 128]); } __syncthreads(); }
    if(blockSize >= 128) { if (tid <  64 && tid +  64 < n) { sdata[tid] = reduceFunc.CU(sdata[tid], sdata[tid +  64]); } __syncthreads(); }
    if(blockSize >=  64) { if (tid <  32 && tid +  32 < n) { sdata[tid] = reduceFunc.CU(sdata[tid], sdata[tid +  32]); } __syncthreads(); }
    if(blockSize >=  32) { if (tid <  16 && tid +  16 < n) { sdata[tid] = reduceFunc.CU(sdata[tid], sdata[tid +  16]); } __syncthreads(); }
    if(blockSize >=  16) { if (tid <   8 && tid +   8 < n) { sdata[tid] = reduceFunc.CU(sdata[tid], sdata[tid +   8]); } __syncthreads(); }
    if(blockSize >=   8) { if (tid <   4 && tid +   4 < n) { sdata[tid] = reduceFunc.CU(sdata[tid], sdata[tid +   4]); } __syncthreads(); }
    if(blockSize >=   4) { if (tid <   2 && tid +   2 < n) { sdata[tid] = reduceFunc.CU(sdata[tid], sdata[tid +   2]); } __syncthreads(); }
    if(blockSize >=   2) { if (tid <   1 && tid +   1 < n) { sdata[tid] = reduceFunc.CU(sdata[tid], sdata[tid +   1]); } __syncthreads(); }

    if(tid == 0)
    {
        output[blockIdx.x] = sdata[tid];
    }
}

}

#endif

#endif