mapoverlap_kernels.h

Go to the documentation of this file.

#ifndef MAPOVERLAP_KERNELS_H
#define MAPOVERLAP_KERNELS_H

#ifdef SKEPU_OPENCL

#include <string>

namespace skepu
{

static std::string MatrixTranspose_CL(
"__kernel void matrix_transpose_KERNELNAME(__global float *odata, __global float *idata, int offset, int width, int height, __local float* block)\n"
"{\n"
"   unsigned int xIndex = get_global_id(0);\n"
"   unsigned int yIndex = get_global_id(1);\n"
"   if((xIndex + offset < width) && (yIndex < height))\n"
"   {\n"
"       unsigned int index_in = yIndex * width + xIndex + offset;\n"
"       block[get_local_id(1)*(BLOCK_DIM+1)+get_local_id(0)] = idata[index_in];\n"
"   }\n"
"   barrier(CLK_LOCAL_MEM_FENCE);\n"
"   xIndex = get_group_id(1) * BLOCK_DIM + get_local_id(0);\n"
"   yIndex = get_group_id(0) * BLOCK_DIM + get_local_id(1);\n"
"   if((xIndex < height) && (yIndex + offset < width))\n"
"       {\n"
"       unsigned int index_out = yIndex * height + xIndex;\n"
"       odata[index_out] = block[get_local_id(0)*(BLOCK_DIM+1)+get_local_id(1)];\n"
"   }\n"
"}\n"
);



static std::string MapOverlapKernel_CL(
"__kernel void MapOverlapKernel_KERNELNAME(__global TYPE* input, __global TYPE* output, __global TYPE* wrap, int n, int overlap, int out_offset, int out_numelements, int poly, TYPE pad, __local TYPE* sdata)\n"
"{\n"
"    int tid = get_local_id(0);\n"
"    int i = get_group_id(0) * get_local_size(0) + get_local_id(0);\n"
"    if(poly == 0)\n"
"    {\n"
"        sdata[overlap+tid] = (i < n) ? input[i] : pad;\n"
"        if(tid < overlap)\n"
"        {\n"
"            sdata[tid] = (get_group_id(0) == 0) ? pad : input[i-overlap];\n"
"        }\n"
"        if(tid >= (get_local_size(0)-overlap))\n"
"        {\n"
"            sdata[tid+2*overlap] = (get_group_id(0) != get_num_groups(0)-1 && i+overlap < n) ? input[i+overlap] : pad;\n"
"        }\n"
"    }\n"
"    else if(poly == 1)\n"
"    {\n"
"        if(i < n)\n"
"        {\n"
"           sdata[overlap+tid] = input[i];\n"
"        }\n"
"        else if(i-n < overlap)\n"
"        {\n"
"           sdata[overlap+tid] = wrap[overlap+(i-n)];\n"
"        }\n"
"        else\n"
"        {\n"
"           sdata[overlap+tid] = pad;\n"
"        }\n"
"        if(tid < overlap)\n"
"        {\n"
"               sdata[tid] = (get_group_id(0) == 0) ? wrap[tid] : input[i-overlap];\n"
"        }\n"
"        if(tid >= (get_local_size(0)-overlap))\n"
"        {\n"
"               sdata[tid+2*overlap] = (get_group_id(0) != get_num_groups(0)-1 && i+overlap < n) ? input[i+overlap] : wrap[overlap+(i+overlap-n)];\n"
"        }\n"
"    }\n"
"    else if(poly == 2)\n"
"    {\n"
"        sdata[overlap+tid] = (i < n) ? input[i] : input[n-1];\n"
"        if(tid < overlap)\n"
"        {\n"
"            sdata[tid] = (get_group_id(0) == 0) ? input[0] : input[i-overlap];\n"
"        }\n"
"        if(tid >= (get_local_size(0)-overlap))\n"
"        {\n"
"            sdata[tid+2*overlap] = (get_group_id(0) != get_num_groups(0)-1 && i+overlap < n) ? input[i+overlap] : input[n-1];\n"
"        }\n"
"    }\n"
"    barrier(CLK_LOCAL_MEM_FENCE);\n"
"    if( (i >= out_offset) && (i < out_offset+out_numelements) )\n"
"    {\n"
"        output[i-out_offset] = FUNCTIONNAME(&(sdata[tid+overlap]));\n"
"    }\n"
"}\n"
);


static std::string MapOverlapKernel_CL_Matrix_Row(
"__kernel void MapOverlapKernel_MatRowWise_KERNELNAME(__global TYPE* input, __global TYPE* output, __global TYPE* wrap, int n, int overlap, int out_offset, int out_numelements, int poly, TYPE pad, int blocksPerRow, int rowWidth, __local TYPE* sdata)\n"
"{\n"
"    int tid = get_local_id(0);\n"
"    int i = get_group_id(0) * get_local_size(0) + get_local_id(0);\n"
"    int wrapIndex= 2 * overlap * (int)(get_group_id(0)/blocksPerRow);\n"
"    int tmp= (get_group_id(0) % blocksPerRow);\n"
"    int tmp2= (get_group_id(0) / blocksPerRow);\n"
"    if(poly == 0)\n"
"    {\n"
"        sdata[overlap+tid] = (i < n) ? input[i] : pad;\n"
"        if(tid < overlap)\n"
"        {\n"
"            sdata[tid] = (tmp==0) ? pad : input[i-overlap];\n"
"        }\n"
"        if(tid >= (get_local_size(0)-overlap))\n"
"        {\n"
"            sdata[tid+2*overlap] = (get_group_id(0) != (get_num_groups(0)-1) && (i+overlap < n) && tmp!=(blocksPerRow-1)) ? input[i+overlap] : pad;\n"
"        }\n"
"    }\n"
"    else if(poly == 1)\n"
"    {\n"
"        if(i < n)\n"
"        {\n"
"           sdata[overlap+tid] = input[i];\n"
"        }\n"
"        else if(i-n < overlap)\n"
"        {\n"
"           sdata[overlap+tid] = wrap[(overlap+(i-n))+ wrapIndex];\n"
"        }\n"
"        else\n"
"        {\n"
"           sdata[overlap+tid] = pad;\n"
"        }\n"
"        if(tid < overlap)\n"
"        {\n"
"               sdata[tid] = (tmp==0) ? wrap[tid+wrapIndex] : input[i-overlap];\n"
"        }\n"
"        if(tid >= (get_local_size(0)-overlap))\n"
"        {\n"
"               sdata[tid+2*overlap] = (get_group_id(0) != (get_num_groups(0)-1) && i+overlap < n && tmp!=(blocksPerRow-1)) ? input[i+overlap] : wrap[overlap+wrapIndex+(tid+overlap-get_local_size(0))];\n"
"        }\n"
"    }\n"
"    else if(poly == 2)\n"
"    {\n"
"        sdata[overlap+tid] = (i < n) ? input[i] : input[n-1];\n"
"        if(tid < overlap)\n"
"        {\n"
"            sdata[tid] = (tmp==0) ? input[tmp2*rowWidth] : input[i-overlap];\n"
"        }\n"
"        if(tid >= (get_local_size(0)-overlap))\n"
"        {\n"
"            sdata[tid+2*overlap] = (get_group_id(0) != (get_num_groups(0)-1) && (i+overlap < n) && (tmp!=(blocksPerRow-1))) ? input[i+overlap] : input[(tmp2+1)*rowWidth-1];\n"
"        }\n"
"    }\n"
"    barrier(CLK_LOCAL_MEM_FENCE);\n"
"    if( (i >= out_offset) && (i < out_offset+out_numelements) )\n"
"    {\n"
"        output[i-out_offset] = FUNCTIONNAME(&(sdata[tid+overlap]));\n"
"    }\n"
"}\n"
);



static std::string MapOverlapKernel_CL_Matrix_Col(
"__kernel void MapOverlapKernel_MatColWise_KERNELNAME(__global TYPE* input, __global TYPE* output, __global TYPE* wrap, int n, int overlap, int out_offset, int out_numelements, int poly, TYPE pad, int blocksPerCol, int rowWidth, int colWidth, __local TYPE* sdata)\n"
"{\n"
"    int tid = get_local_id(0);\n"
"    int i = get_group_id(0) * get_local_size(0) + get_local_id(0);\n"
"    int wrapIndex= 2 * overlap * (int)(get_group_id(0)/blocksPerCol);\n"
"    int tmp= (get_group_id(0) % blocksPerCol);\n"
"    int tmp2= (get_group_id(0) / blocksPerCol);\n"
"    int arrInd = (tid + tmp*get_local_size(0))*rowWidth + tmp2;\n"
"    if(poly == 0)\n"
"    {\n"
"        sdata[overlap+tid] = (i < n) ? input[arrInd] : pad;\n"
"        if(tid < overlap)\n"
"        {\n"
"            sdata[tid] = (tmp==0) ? pad : input[(arrInd-(overlap*rowWidth))];\n"
"        }\n"
"        if(tid >= (get_local_size(0)-overlap))\n"
"        {\n"
"            sdata[tid+2*overlap] = (get_group_id(0) != (get_num_groups(0)-1) && (arrInd+(overlap*rowWidth)) < n && (tmp!=(blocksPerCol-1))) ? input[(arrInd+(overlap*rowWidth))] : pad;\n"
"        }\n"
"    }\n"
"    else if(poly == 1)\n"
"    {\n"
"        if(i < n)\n"
"        {\n"
"           sdata[overlap+tid] = input[arrInd];\n"
"        }\n"
"        else if(i-n < overlap)\n"
"        {\n"
"           sdata[overlap+tid] = wrap[(overlap+(i-n))+ wrapIndex];\n"
"        }\n"
"        else\n"
"        {\n"
"           sdata[overlap+tid] = pad;\n"
"        }\n"
"        if(tid < overlap)\n"
"        {\n"
"               sdata[tid] = (tmp==0) ? wrap[tid+wrapIndex] : input[(arrInd-(overlap*rowWidth))];\n"
"        }\n"
"        if(tid >= (get_local_size(0)-overlap))\n"
"        {\n"
"               sdata[tid+2*overlap] = (get_group_id(0) != (get_num_groups(0)-1) && (arrInd+(overlap*rowWidth)) < n && (tmp!=(blocksPerCol-1))) ? input[(arrInd+(overlap*rowWidth))] : wrap[overlap+wrapIndex+(tid+overlap-get_local_size(0))];\n"
"        }\n"
"    }\n"
"    else if(poly == 2)\n"
"    {\n"
"        sdata[overlap+tid] = (i < n) ? input[arrInd] : input[n-1];\n"
"        if(tid < overlap)\n"
"        {\n"
"            sdata[tid] = (tmp==0) ? input[tmp2] : input[(arrInd-(overlap*rowWidth))];\n"
"        }\n"
"        if(tid >= (get_local_size(0)-overlap))\n"
"        {\n"
"            sdata[tid+2*overlap] = (get_group_id(0) != (get_num_groups(0)-1) && (arrInd+(overlap*rowWidth)) < n && (tmp!=(blocksPerCol-1))) ? input[(arrInd+(overlap*rowWidth))] : input[tmp2+(colWidth-1)*rowWidth];\n"
"        }\n"
"    }\n"
"    barrier(CLK_LOCAL_MEM_FENCE);\n"
"    if( (i >= out_offset) && (i < out_offset+out_numelements) )\n"
"    {\n"
"        output[arrInd-out_offset] = FUNCTIONNAME(&(sdata[tid+overlap]));\n"
"    }\n"
"}\n"
);













}

#endif

#ifdef SKEPU_CUDA

namespace skepu
{

// #define BLOCK_DIM 16

template <typename T>
__global__ void transpose(T *odata, T *idata, int width, int height)
{
//  __shared__ T block[BLOCK_DIM][BLOCK_DIM+1];
        extern __shared__ char _sdata[];
    T* sdata = reinterpret_cast<T*>(_sdata);

    int block_dim= blockDim.x;
    int block_dimY= blockDim.y;
    // read the matrix tile into shared memory
    unsigned int xIndex = blockIdx.x * block_dim + threadIdx.x;
    unsigned int yIndex = blockIdx.y * block_dimY + threadIdx.y;
    if((xIndex < width) && (yIndex < height))
    {
        unsigned int index_in = yIndex * width + xIndex;
        sdata[threadIdx.y][threadIdx.x] = idata[index_in];
    }

        __syncthreads();

    // write the transposed matrix tile to global memory
    xIndex = blockIdx.y * block_dim + threadIdx.x;
    yIndex = blockIdx.x * block_dimY + threadIdx.y;
    if((xIndex < height) && (yIndex < width))
    {
        unsigned int index_out = yIndex * height + xIndex;
        odata[index_out] = sdata[threadIdx.x][threadIdx.y];
    }
}

template <int poly, typename T, typename OverlapFunc>
__global__ void MapOverlapKernel_CU(OverlapFunc mapOverlapFunc, T* input, T* output, T* wrap, unsigned int n, unsigned int out_offset, unsigned int out_numelements, T pad)
{
    extern __shared__ char _sdata[];
    T* sdata = reinterpret_cast<T*>(_sdata);

    unsigned int tid = threadIdx.x;
    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
    int overlap = mapOverlapFunc.overlap;

    //Copy data to shared memory
    if(poly == 0)
    {
        sdata[overlap+tid] = (i < n) ? input[i] : pad;

        if(tid < overlap)
        {
            sdata[tid] = (blockIdx.x == 0) ? pad : input[i-overlap];
        }

        if(tid >= (blockDim.x-overlap))
        {
            sdata[tid+2*overlap] = (blockIdx.x != gridDim.x-1 && i+overlap < n) ? input[i+overlap] : pad;
        }
    }
    else if(poly == 1)
    {
        if(i < n)
        {
            sdata[overlap+tid] = input[i];
        }
        else if(i-n < overlap)
        {
            sdata[overlap+tid] = wrap[overlap+(i-n)];
        }
        else
        {
            sdata[overlap+tid] = pad;
        }

        if(tid < overlap)
        {
            sdata[tid] = (blockIdx.x == 0) ? wrap[tid] : input[i-overlap];
        }

        if(tid >= (blockDim.x-overlap))
        {
            sdata[tid+2*overlap] = (blockIdx.x != gridDim.x-1 && i+overlap < n) ? input[i+overlap] : wrap[overlap+(i+overlap-n)];
        }
    }
    else if(poly == 2)
    {
        sdata[overlap+tid] = (i < n) ? input[i] : input[n-1];

        if(tid < overlap)
        {
            sdata[tid] = (blockIdx.x == 0) ? input[0] : input[i-overlap];
        }

        if(tid >= (blockDim.x-overlap))
        {
            sdata[tid+2*overlap] = (blockIdx.x != gridDim.x-1 && i+overlap < n) ? input[i+overlap] : input[n-1];
        }
    }

    __syncthreads();

    //Compute and store data
    if( (i >= out_offset) && (i < out_offset+out_numelements) )
    {
        output[i-out_offset] = mapOverlapFunc.CU(&(sdata[tid+overlap]));
    }
}






template <int poly, typename T, typename OverlapFunc>
__global__ void MapOverlapKernel_CU_Matrix_Row(OverlapFunc mapOverlapFunc, T* input, T* output, T* wrap, unsigned int n, unsigned int out_offset, unsigned int out_numelements, T pad, unsigned int blocksPerRow, unsigned int rowWidth)
{
    extern __shared__ char _sdata[];
    T* sdata = reinterpret_cast<T*>(_sdata);

    unsigned int tid = threadIdx.x;
    unsigned int i = blockIdx.x * blockDim.x + tid;
    int overlap = mapOverlapFunc.overlap;
    
    unsigned wrapIndex= 2 * overlap * (int)(blockIdx.x/blocksPerRow);
    int tmp= (blockIdx.x % blocksPerRow);
    int tmp2= (blockIdx.x / blocksPerRow);

    
    //Copy data to shared memory
    if(poly == 0)
    {
        sdata[overlap+tid] = (i < n) ? input[i] : pad;

        if(tid < overlap)
        {
            sdata[tid] = (tmp==0) ? pad : input[i-overlap];
        }

        if(tid >= (blockDim.x-overlap))
        {
            sdata[tid+2*overlap] = (blockIdx.x != gridDim.x-1 && (i+overlap < n) && tmp!=(blocksPerRow-1)) ? input[i+overlap] : pad;
        }
    }
    else if(poly == 1)
    {
        if(i < n)
        {
            sdata[overlap+tid] = input[i];
        }
        else if(i-n < overlap)
        {
            sdata[overlap+tid] = wrap[(overlap+(i-n))+ wrapIndex];
        }
        else
        {
            sdata[overlap+tid] = pad;
        }

        if(tid < overlap)
        {
            sdata[tid] = (tmp==0) ? wrap[tid+wrapIndex] : input[i-overlap];
        }

        if(tid >= (blockDim.x-overlap))
        {
        sdata[tid+2*overlap] = (blockIdx.x != gridDim.x-1 && i+overlap < n && tmp!=(blocksPerRow-1)) ? input[i+overlap] : wrap[overlap+wrapIndex+(tid+overlap-blockDim.x)];
        }
    }
    else if(poly == 2)
    {
        sdata[overlap+tid] = (i < n) ? input[i] : input[n-1];

        if(tid < overlap)
        {
            sdata[tid] = (tmp==0) ? input[tmp2*rowWidth] : input[i-overlap];
        }

        if(tid >= (blockDim.x-overlap))
        {
            sdata[tid+2*overlap] = (blockIdx.x != gridDim.x-1 && (i+overlap < n) && (tmp!=(blocksPerRow-1))) ? input[i+overlap] : input[(tmp2+1)*rowWidth-1];
        }
    }

    __syncthreads();

    //Compute and store data
    if( (i >= out_offset) && (i < out_offset+out_numelements) )
    {
        output[i-out_offset] = mapOverlapFunc.CU(&(sdata[tid+overlap]));
    }
}






template <int poly, typename T, typename OverlapFunc>
__global__ void MapOverlapKernel_CU_Matrix_Col(OverlapFunc mapOverlapFunc, T* input, T* output, T* wrap, unsigned int n, unsigned int out_offset, unsigned int out_numelements, T pad, unsigned int blocksPerCol, unsigned int rowWidth, unsigned int colWidth)
{
    extern __shared__ char _sdata[];
    T* sdata = reinterpret_cast<T*>(_sdata);

    unsigned int tid = threadIdx.x;
    unsigned int i = blockIdx.x * blockDim.x + tid;
    int overlap = mapOverlapFunc.overlap;
    
    unsigned wrapIndex= 2 * overlap * (int)(blockIdx.x/blocksPerCol);
    int tmp= (blockIdx.x % blocksPerCol);
    int tmp2= (blockIdx.x / blocksPerCol);

    unsigned int arrInd = (threadIdx.x + tmp*blockDim.x)*rowWidth + ((blockIdx.x)/blocksPerCol);
    
    //Copy data to shared memory
    if(poly == 0)
    {
        sdata[overlap+tid] = (i < n) ? input[arrInd] : pad;

        if(tid < overlap)
        {
            sdata[tid] = (tmp==0) ? pad : input[(arrInd-(overlap*rowWidth))];
        }

        if(tid >= (blockDim.x-overlap))
        {
            sdata[tid+2*overlap] = (blockIdx.x != gridDim.x-1 && (arrInd+(overlap*rowWidth)) < n && (tmp!=(blocksPerCol-1))) ? input[(arrInd+(overlap*rowWidth))] : pad;
        }
    }
    else if(poly == 1)
    {
        if(i < n)
        {
            sdata[overlap+tid] = input[arrInd];
        }
        else if(i-n < overlap)
        {
            sdata[overlap+tid] = wrap[(overlap+(i-n))+ wrapIndex];
        }
        else
        {
            sdata[overlap+tid] = pad;
        }

        if(tid < overlap)
        {
            sdata[tid] = (tmp==0) ? wrap[tid+wrapIndex] : input[(arrInd-(overlap*rowWidth))];
        }

        if(tid >= (blockDim.x-overlap))
        {
            sdata[tid+2*overlap] = (blockIdx.x != gridDim.x-1 && (arrInd+(overlap*rowWidth)) < n && (tmp!=(blocksPerCol-1))) ? input[(arrInd+(overlap*rowWidth))] : wrap[overlap+wrapIndex+(tid+overlap-blockDim.x)];
        }
    }
    else if(poly == 2)
    {
        sdata[overlap+tid] = (i < n) ? input[arrInd] : input[n-1];

        if(tid < overlap)
        {
            sdata[tid] = (tmp==0) ? input[tmp2] : input[(arrInd-(overlap*rowWidth))];
        }

        if(tid >= (blockDim.x-overlap))
        {
            sdata[tid+2*overlap] = (blockIdx.x != gridDim.x-1 && (arrInd+(overlap*rowWidth)) < n && (tmp!=(blocksPerCol-1))) ? input[(arrInd+(overlap*rowWidth))] : input[tmp2+(colWidth-1)*rowWidth];
        }
    }

    __syncthreads();

    //Compute and store data
    if( (arrInd >= out_offset) && (arrInd < out_offset+out_numelements) )
    {
        output[arrInd-out_offset] = mapOverlapFunc.CU(&(sdata[tid+overlap]));
    }
}






}

#endif

#endif