maparray_kernels.h

Go to the documentation of this file.

#ifndef MAPARRAY_KERNELS_H
#define MAPARRAY_KERNELS_H

#ifdef SKEPU_OPENCL

#include <string>

namespace skepu
{

static std::string MapArrayKernel_CL(
"__kernel void MapArrayKernel_KERNELNAME(__global TYPE* input1, __global TYPE* input2, __global TYPE* output, unsigned 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"
"        output[i] = FUNCTIONNAME(&input1[0], input2[i]);\n"
"        i += gridSize;\n"
"    }\n"
"}\n"
);


//static std::string MapArrayKernel_CL_Matrix(
//"__kernel void MapArrayKernel_KERNELNAME_Matrix(__global TYPE* input1, __global TYPE* input2, __global TYPE* output, unsigned int n, unsigned int xsize, unsigned int ysize)\n"
//"{\n"
//"    int xindex = get_global_id(0);\n"
//"    int yindex = get_global_id(1);\n"
//"    int i = (get_global_size(0) * get_local_size(0)) * yindex + xindex;\n"
//"    if(i < n && xindex<xsize && yindex <ysize)\n"
//"    {\n"
//"        output[i] = FUNCTIONNAME(&input1[0], input2[i]);\n"
//"    }\n"
//"}\n"
//);

static std::string MapArrayKernel_CL_Matrix(
"__kernel void MapArrayKernel_Matrix_KERNELNAME(__global TYPE* input1, __global TYPE* input2, __global TYPE* output, unsigned int n, unsigned int xsize, unsigned int ysize, int yoffset, CONST_TYPE const1)\n"
"{\n"
"    int xindex = get_global_id(0);\n"
"    int yindex = get_global_id(1);\n"
"    int i = yindex*xsize + xindex; \n"
"    if(i < n && xindex<xsize && yindex <ysize)\n"
"    {\n"
"        output[i] = FUNCTIONNAME(&input1[0], input2[i], xindex, yindex+yoffset, const1);\n"
"    }\n"
"}\n"
);


static std::string MapArrayKernel_CL_Matrix_Blockwise(
"__kernel void MapArrayKernel_Matrix_Blockwise_KERNELNAME(__global TYPE* input1, __global TYPE* input2, __global TYPE* output, unsigned int outSize, int p2BlockSize, CONST_TYPE const1)\n"
"{\n"
"    int i = get_global_id(0);\n"
"    unsigned int gridSize = get_local_size(0)*get_num_groups(0);\n"
"    if(i < outSize)\n"
"    {\n"
"        output[i] = FUNCTIONNAME(&input1[0], &input2[i*p2BlockSize], const1);\n"
"        i += gridSize;\n"
"    }\n"
"}\n"
);


static std::string MapArrayKernel_CL_Sparse_Matrix_Blockwise(
"__kernel void MapArrayKernel_Sparse_Matrix_Blockwise_KERNELNAME(__global TYPE* input1, __global TYPE* in2_values, __global unsigned int *in2_row_offsets, __global unsigned int *in2_col_indices, __global TYPE* output, unsigned int outSize, int indexOffset, CONST_TYPE const1)\n"
"{\n"
"    int i = get_global_id(0);\n"
"    unsigned int gridSize = get_local_size(0)*get_num_groups(0);\n"
"    if(i < outSize)\n"
"    {\n"
"        int rowId = in2_row_offsets[i] - indexOffset;\n"
"        int row2Id = in2_row_offsets[i+1] - indexOffset;\n"
"        output[i] = FUNCTIONNAME(&input1[0], &in2_values[rowId], (row2Id-rowId), &in2_col_indices[rowId], const1);\n"
"        i += gridSize;\n"
"    }\n"
"}\n"
);

}

#endif

#ifdef SKEPU_CUDA

namespace skepu
{

template <typename T, typename ArrayFunc>
__global__ void MapArrayKernel_CU(ArrayFunc mapArrayFunc, T* input1, T* input2, T* output, unsigned int n)
{
    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int gridSize = blockDim.x*gridDim.x;

    while(i < n)
    {
        output[i] = mapArrayFunc.CU(&input1[0], input2[i]);
        i += gridSize;
    }
}

//
//
//template <typename T, typename ArrayFunc>
//__global__ void MapArrayKernel_CU_Matrix(ArrayFunc mapArrayFunc, T* input1, T* input2, T* output, unsigned int n)
//{
//  unsigned int xindex  = blockIdx.x * blockDim.x + threadIdx.x;
//    unsigned int yindex  = blockIdx.y * blockDim.y + threadIdx.y;
//    unsigned int outaddr = (gridDim.x * blockDim.x) * yindex + xindex;
//
//    if(outaddr < n)
//    {
//        output[outaddr] = mapArrayFunc.CU(&input1[0], input2[outaddr], xindex, yindex);
//    }
//}

template <typename T, typename ArrayFunc>
__global__ void MapArrayKernel_CU_Matrix(ArrayFunc mapArrayFunc, T* input1, T* input2, T* output, unsigned int n, unsigned int xsize, unsigned int ysize, unsigned int yoffset)
{
    unsigned int xindex  = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int yindex  = blockIdx.y * blockDim.y + threadIdx.y;
    unsigned int outaddr = yindex*xsize + xindex; //(gridDim.x * blockDim.x) * yindex + xindex;

    if(outaddr < n && xindex<xsize && yindex <ysize)
    {
        output[outaddr] = mapArrayFunc.CU(&input1[0], input2[outaddr], xindex, yindex+yoffset);
    }
}


template <typename T, typename ArrayFunc>
__global__ void MapArrayKernel_CU_Matrix_Blockwise(ArrayFunc mapArrayFunc, T* input1, T* input2, T* output, unsigned int outSize, int p2BlockSize)
{
    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int gridSize = blockDim.x*gridDim.x;

    if(i < outSize)
    {
        output[i] = mapArrayFunc.CU(&input1[0], &input2[i*p2BlockSize]);
        i += gridSize;
    }
}


template <typename T, typename ArrayFunc>
__global__ void MapArrayKernel_CU_Sparse_Matrix_Blockwise(ArrayFunc mapArrayFunc, T* input1, T* in2_values, unsigned int *in2_row_offsets, unsigned int *in2_col_indices, T* output, unsigned int outSize, int indexOffset)
{
    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int gridSize = blockDim.x*gridDim.x;

    if(i < outSize)    //m_mapArrayFunc->CPU(&input1[0], &values[rowIdx], rowSize, &col_indices[rowIdx]);
    {  
        int rowId = in2_row_offsets[i] - indexOffset;
        int row2Id = in2_row_offsets[i+1] - indexOffset;
        output[i] = mapArrayFunc.CU(&input1[0], &in2_values[rowId], (row2Id-rowId), &in2_col_indices[rowId]);
        i += gridSize;
    }
}





}

#endif

#endif