include/skepu/src/scan_kernels.h

Go to the documentation of this file.

#ifndef SCAN_KERNELS_H
#define SCAN_KERNELS_H

#ifdef SKEPU_OPENCL

#include <string>

namespace skepu
{

static std::string ScanKernel_CL(
"__kernel void ScanKernel_KERNELNAME(__global TYPE* input, __global TYPE* output, __global TYPE* blockSums, unsigned int n, unsigned int numElements, __local TYPE* sdata)\n"
"{\n"
"    unsigned int threadIdx = get_local_id(0);\n"
"    unsigned int blockDim = get_local_size(0);\n"
"    unsigned int blockIdx = get_group_id(0);\n"
"    unsigned int gridDim = get_num_groups(0);\n"
"    int thid = threadIdx;\n"
"    int pout = 0;\n"
"    int pin = 1;\n"
"    int mem = get_global_id(0);\n"
"    int blockNr = blockIdx;\n"
"    unsigned int gridSize = blockDim*gridDim;\n"
"    unsigned int numBlocks = numElements/(blockDim) + (numElements%(blockDim) == 0 ? 0:1);\n"
"    int offset;\n"
"    while(blockNr < numBlocks)\n"
"    {\n"
"        sdata[pout*n+thid] = (mem < numElements) ? input[mem] : 0;\n"
"        barrier(CLK_LOCAL_MEM_FENCE);\n"
"        for(offset = 1; offset < n; offset *=2)\n"
"        {\n"
"            pout = 1-pout;\n"
"            pin = 1-pout;\n"
"            if(thid >= offset)\n"
"                sdata[pout*n+thid] = FUNCTIONNAME(sdata[pin*n+thid], sdata[pin*n+thid-offset], (TYPE)0);\n"
"            else\n"
"                sdata[pout*n+thid] = sdata[pin*n+thid];\n"
"            barrier(CLK_LOCAL_MEM_FENCE);\n"
"        }\n"
"        if(thid == blockDim - 1)\n"
"            blockSums[blockNr] = sdata[pout*n+blockDim-1];\n"
"        if(mem < numElements)\n"
"            output[mem] = sdata[pout*n+thid];\n"
"        mem += gridSize;\n"
"        blockNr += gridDim;\n"
"        barrier(CLK_LOCAL_MEM_FENCE);\n"
"    }\n"
"}\n"
);

static std::string ScanUpdate_CL(
"__kernel void ScanUpdate_KERNELNAME(__global TYPE* data, __global TYPE* sums, int isInclusive, TYPE init, int n, __global TYPE* ret, __local TYPE* sdata)\n"
"{\n"
"    __local TYPE offset;\n"
"    __local TYPE inc_offset;\n"
"    unsigned int threadIdx = get_local_id(0);\n"
"    unsigned int blockDim = get_local_size(0);\n"
"    unsigned int blockIdx = get_group_id(0);\n"
"    unsigned int gridDim = get_num_groups(0);\n"
"    int thid = threadIdx;\n"
"    int blockNr = blockIdx;\n"
"    unsigned int gridSize = blockDim*gridDim;\n"
"    int mem = get_global_id(0);\n"
"    unsigned int numBlocks = n/(blockDim) + (n%(blockDim) == 0 ? 0:1);\n"
"    while(blockNr < numBlocks)\n"
"    {\n"
"        if(thid == 0)\n"
"        {\n"
"            if(isInclusive == 0)\n"
"            {\n"
"                offset = init;\n"
"                if(blockNr > 0)\n"
"                {\n"
"                    offset = FUNCTIONNAME(offset, sums[blockNr-1], (TYPE)0);\n"
"                    inc_offset = sums[blockNr-1];\n"
"                }\n"
"            }\n"
"            else\n"
"            {\n"
"                if(blockNr > 0)\n"
"                    offset = sums[blockNr-1];\n"
"            }\n"
"        }\n"
"        barrier(CLK_LOCAL_MEM_FENCE);\n"
"        if(isInclusive == 1)\n"
"        {\n"
"            if(blockNr > 0)\n"
"                sdata[thid] = (mem < n) ? FUNCTIONNAME(offset, data[mem], (TYPE)0) : 0;\n"
"            else\n"
"                sdata[thid] = (mem < n) ? data[mem] : 0;\n"
"            if(mem == n-1)\n"
"                *ret = sdata[thid];\n"
"        }\n"
"        else\n"
"        {\n"
"            if(mem == n-1)\n"
"                *ret = FUNCTIONNAME(inc_offset, data[mem], (TYPE)0);\n"
"            if(thid == 0)\n"
"                sdata[thid] = offset;\n"
"            else\n"
"                sdata[thid] = (mem-1 < n) ? FUNCTIONNAME(offset, data[mem-1], (TYPE)0) : 0;\n"
"        }\n"
"        barrier(CLK_LOCAL_MEM_FENCE);\n"
"        if(mem < n)\n"
"            data[mem] = sdata[thid];\n"
"        mem += gridSize;\n"
"        blockNr += gridDim;\n"
"        barrier(CLK_LOCAL_MEM_FENCE);\n"
"    }\n"
"}\n"
);

static std::string ScanAdd_CL(
"__kernel void ScanAdd_KERNELNAME(__global TYPE* data, TYPE sum, int n)\n"
"{\n"
"    int i = get_global_id(0);\n"
"    unsigned int gridSize = get_local_size(0)*get_num_groups(0);\n"
"    while(i < n)\n"
"    {\n"
"        data[i] = FUNCTIONNAME(data[i], sum, (TYPE)0);\n"
"        i += gridSize;\n"
"    }\n"
"}\n"
);

}

#endif

#ifdef SKEPU_CUDA

namespace skepu
{

template<typename T, typename BinaryFunc>
__global__ void ScanKernel_CU(BinaryFunc scanFunc, T* input, T* output, T* blockSums, unsigned int n, unsigned int numElements)
{
    extern __shared__ char _sdata[];
    T* sdata = reinterpret_cast<T*>(_sdata);

    int thid = threadIdx.x;
    int pout = 0;
    int pin = 1;
    int mem = __mul24(blockIdx.x, blockDim.x) + threadIdx.x;
    int blockNr = blockIdx.x;
    unsigned int gridSize = blockDim.x*gridDim.x;
    unsigned int numBlocks = numElements/(blockDim.x) + (numElements%(blockDim.x) == 0 ? 0:1);
    int offset;

    while(blockNr < numBlocks)
    {
        sdata[pout*n+thid] = (mem < numElements) ? input[mem] : 0;

        __syncthreads();

        for(offset = 1; offset < n; offset *=2)
        {
            pout = 1-pout;
            pin = 1-pout;
            if(thid >= offset)
                sdata[pout*n+thid] = scanFunc.CU(sdata[pin*n+thid], sdata[pin*n+thid-offset]);
            else
                sdata[pout*n+thid] = sdata[pin*n+thid];

            __syncthreads();
        }
        if(thid == blockDim.x - 1)
            blockSums[blockNr] = sdata[pout*n+blockDim.x-1];

        if(mem < numElements)
            output[mem] = sdata[pout*n+thid];

        mem += gridSize;
        blockNr += gridDim.x;

        __syncthreads();
    }
}

template <typename T, typename BinaryFunc>
__global__ void ScanUpdate_CU(BinaryFunc scanFunc, T *data, T *sums, int isInclusive, T init, int n, T *ret)
{
    extern __shared__ char _sdata[];
    T* sdata = reinterpret_cast<T*>(_sdata);

    __shared__ T offset;
    __shared__ T inc_offset;

    int thid = threadIdx.x;
    int blockNr = blockIdx.x;
    unsigned int gridSize = blockDim.x*gridDim.x;
    int mem = __mul24(blockIdx.x, blockDim.x) + threadIdx.x;
    unsigned int numBlocks = n/(blockDim.x) + (n%(blockDim.x) == 0 ? 0:1);

    while(blockNr < numBlocks)
    {
        if(thid == 0)
        {
            if(isInclusive == 0)
            {
                offset = init;
                if(blockNr > 0)
                {
                    offset = scanFunc.CU(offset, sums[blockNr-1]);
                    inc_offset = sums[blockNr-1];
                }
            }
            else
            {
                if(blockNr > 0)
                    offset = sums[blockNr-1];
            }
        }

        __syncthreads();

        if(isInclusive == 1)
        {
            if(blockNr > 0)
                sdata[thid] = (mem < n) ? scanFunc.CU(offset, data[mem]) : 0;
            else
                sdata[thid] = (mem < n) ? data[mem] : 0;
            if(mem == n-1)
                *ret = sdata[thid];
        }
        else
        {
            if(mem == n-1)
                *ret = scanFunc.CU(inc_offset, data[mem]);
            if(thid == 0)
                sdata[thid] = offset;
            else
                sdata[thid] = (mem-1 < n) ? scanFunc.CU(offset, data[mem-1]) : 0;
        }

        __syncthreads();

        if(mem < n)
            data[mem] = sdata[thid];

        mem += gridSize;
        blockNr += gridDim.x;

        __syncthreads();
    }
}

template <typename T, typename BinaryFunc>
__global__ void ScanAdd_CU(BinaryFunc scanFunc, T *data, T sum, int n)
{
    int i = __mul24(blockIdx.x, blockDim.x) + threadIdx.x;
    unsigned int gridSize = blockDim.x*gridDim.x;

    while(i < n)
    {
        data[i] = scanFunc.CU(data[i], sum);
        i += gridSize;
    }
}

}

#endif

#endif