generate_kernels.h

Go to the documentation of this file.

#ifndef GENERATE_KERNELS_H
#define GENERATE_KERNELS_H

#ifdef SKEPU_OPENCL

#include <string>

namespace skepu
{

static std::string GenerateKernel_CL(
"__kernel void GenerateKernel_KERNELNAME(__global TYPE* output, unsigned int numElements, unsigned int offset, CONST_TYPE const1)\n"
"{\n"
"    output = (__global void *)output + offset; /* partitioning is special with opencl */ \n"
"    unsigned int i = get_global_id(0);\n"
"    unsigned int gridSize = get_local_size(0)*get_num_groups(0);\n"
"    while(i < numElements)\n"
"    {\n"
"        output[i] = FUNCTIONNAME(i+offset, const1);\n"
"        i += gridSize;\n"
"    }\n"
"}\n"
);


static std::string GenerateKernel_CL_Matrix(
"__kernel void GenerateKernel_Matrix_KERNELNAME(__global TYPE* output, unsigned int numElements, unsigned int xsize, unsigned int ysize, unsigned int offset, CONST_TYPE const1)\n"
"{\n"
"    output = (__global void *)output + offset; /* partitioning is special with opencl */ \n"
"    int xindex = get_global_id(0);\n"
"    int yindex = get_global_id(1);\n"
"    int i = yindex*xsize + xindex; \n"
"    if(i < numElements && xindex<xsize && yindex <ysize)\n"
"    {\n"
"        output[i] = FUNCTIONNAME(xindex, yindex+offset, const1);\n"
"    }\n"
"}\n"
);

}

#endif

#ifdef SKEPU_CUDA

namespace skepu
{

template <typename T, typename GenerateFunc>
__global__ void GenerateKernel_CU(GenerateFunc generateFunc, T* output, unsigned int numElements, unsigned int indexOffset)
{
    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int gridSize = blockDim.x*gridDim.x;

    while(i < numElements)
    {
        output[i] = generateFunc.CU(i+indexOffset);
        i += gridSize;
    }
}



template <typename T, typename GenerateFunc>
__global__ void GenerateKernel_CU_Matrix(GenerateFunc generateFunc, T* output, unsigned int numElements, 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; 

    unsigned int gridSize = blockDim.x*blockDim.y*gridDim.x*gridDim.y;
    while(outaddr < numElements && xindex<xsize && yindex <ysize)
    {
        output[outaddr] = generateFunc.CU(xindex, yindex+yoffset);
        outaddr += gridSize;
        xindex += blockDim.x*gridDim.x;
        yindex += blockDim.y*gridDim.y;
    }
//    if(outaddr < numElements && xindex<xsize && yindex <ysize)
//    {
//        output[outaddr] = generateFunc.CU(xindex, yindex+yoffset);
//    }
    
    

}

}

#endif

#endif