map_kernels.h

Go to the documentation of this file.

#ifndef MAP_KERNELS_H
#define MAP_KERNELS_H

#ifdef SKEPU_OPENCL

#include <string>

namespace skepu
{

static std::string UnaryMapKernel_CL(
"__kernel void UnaryMapKernel_KERNELNAME(__global TYPE* input, __global TYPE* output, unsigned int numElements, unsigned int offset1, unsigned int offset2, CONST_TYPE const1)\n"
"{\n"
"    input = (__global void *)input + offset1; /* partitioning is special with opencl */ \n"
"    output = (__global void *)output + offset2; /* partitioning is special with opencl */ \n"
"    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(input[i], const1);\n"
"        i += gridSize;\n"
"    }\n"
"}\n"
);

static std::string BinaryMapKernel_CL(
"__kernel void BinaryMapKernel_KERNELNAME(__global TYPE* input1, __global TYPE* input2, __global TYPE* output, unsigned int n, unsigned int offset1, unsigned int offset2, unsigned int offset3, CONST_TYPE const1)\n"
"{\n"
"    input1 = (__global void *)input1 + offset1; /* partitioning is special with opencl */ \n"
"    input2 = (__global void *)input2 + offset2; /* partitioning is special with opencl */ \n"
"    output = (__global void *)output + offset3; /* partitioning is special with opencl */ \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[i], input2[i], const1);\n"
"        i += gridSize;\n"
"    }\n"
"}\n"
);

static std::string TrinaryMapKernel_CL(
"__kernel void TrinaryMapKernel_KERNELNAME(__global TYPE* input1, __global TYPE* input2, __global TYPE* input3, __global TYPE* output, unsigned int n, unsigned int offset1, unsigned int offset2, unsigned int offset3, unsigned int offset4, CONST_TYPE const1)\n"
"{\n"
"    input1 = (__global void *)input1 + offset1; /* partitioning is special with opencl */ \n"
"    input2 = (__global void *)input2 + offset2; /* partitioning is special with opencl */ \n"
"    input3 = (__global void *)input3 + offset3; /* partitioning is special with opencl */ \n"
"    output = (__global void *)output + offset4; /* partitioning is special with opencl */ \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[i], input2[i], input3[i], const1);\n"
"        i += gridSize;\n"
"    }\n"
"}\n"
);

}

#endif

#ifdef SKEPU_CUDA

namespace skepu
{

template <typename T, typename U, typename UnaryFunc>
__global__ void MapKernelUnary_CU(UnaryFunc mapFunc, T* input, T* output, unsigned int n, U const1)
{
    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int gridSize = blockDim.x*gridDim.x;

    while(i < n)
    {
        output[i] = mapFunc.CU(input[i], const1);
        i += gridSize;
    }
}

template <typename T, typename U, typename BinaryFunc>
__global__ void MapKernelBinary_CU(BinaryFunc mapFunc, T* input1, T* input2, T* output, unsigned int n, U const1)
{
    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int gridSize = blockDim.x*gridDim.x;

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

template <typename T, typename U, typename TrinaryFunc>
__global__ void MapKernelTrinary_CU(TrinaryFunc mapFunc, T* input1, T* input2, T* input3, T* output, unsigned int n, U const1)
{
    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int gridSize = blockDim.x*gridDim.x;

    while(i < n)
    {
        output[i] = mapFunc.CU(input1[i], input2[i], input3[i], const1);
        i += gridSize;
    }
}

}

#endif

#endif