|
static std::string | skepu::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 | skepu::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 | skepu::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") |
|
Definitions of CUDA and OpenCL kernels for the Map skeleton.
static std::string skepu::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 |
OpenCL Map kernel for binary user functions.
static std::string skepu::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" |
| ) |
|
|
static |
OpenCL Map kernel for trinary user functions.
static std::string skepu::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 |
OpenCL Map kernel for unary user functions.