Opencl sample code

// g++  -g  -o 1 cl-sample.cpp         -lOpenCL -IMali_OpenCL_SDK_v1.1.0/include   -std=c++11

//http://gameeer.com/2015/01/09/opencl-env/

#include <stdio.h>
#include <stdlib.h>
#include <math.h>
#include <CL/opencl.h>
#include <iostream>
// OpenCL kernel. Each work item takes care of one element of c
/* const char *kernelSource =                                          "\n" \  
    "__kernel void vecAdd(  __global float *a,           \n" \
    "                       __global float *b,                       \n" \
    "                       __global float *c,                       \n" \
    "                       const unsigned int n)                    \n" \
    "{                                                               \n" \
    "    //Get our global thread ID                                  \n" \
    "    int id = get_global_id(0);                                  \n" \
    "                                                                \n" \
    "    //Make sure we do not go out of bounds                      \n" \
    "    if (id < n)                                                 \n" \
    "        c[id] = a[id] + b[id];                                  \n" \
    "}                                                               \n" \
    "\n" ;


*/

// OpenCL kernel. Each work item takes care of one element of c
const char *kernelSource =                                          "\n" \
    "__kernel void vecMul(  __global int *a,           \n" \
    "                       __global int *b,                       \n" \
    "                       __global int *c,                       \n" \
    "                       const unsigned int n)                    \n" \
    "{                                                               \n" \
    "    //Get our global thread ID                                  \n" \
    "    int id = get_global_id(0);                                  \n" \
    "                                                                \n" \
    "    //Make sure we do not go out of bounds                      \n" \
    "    if (id < n)                                                 \n" \
    "        c[id] = a[id] * b[id];                                  \n" \
    "}                                                               \n" \
    "\n" ;



   
// Use a static data size for simplicity   
//   
#define IMAGE_X_PIXELS          176   
#define IMAGE_Y_PIXELS          144   
#define IMAGE_SOURCE1_LIMPID    0.5f   
#define IMAGE_SOURCE2_LIMPID    0.5f   

   

// Simple compute kernel which computes the square of an input array    
//   
const char *KernelSource1 = "\n" \
"#define IMAGE_Y_PIXELS   144                                           \n" \
"#define IMAGE_SOURCE1_LIMPID   0.5f                                    \n" \
"#define IMAGE_SOURCE2_LIMPID   0.5f                                    \n" \
"                                                                       \n" \
"__kernel void Limpid(                                                  \n" \
"   __global float image1[][IMAGE_Y_PIXELS],                            \n" \
"   __global float image2[][IMAGE_Y_PIXELS],                            \n" \
"   __global float output[][IMAGE_Y_PIXELS])                            \n" \
"{                                                                      \n" \
"   int x = get_global_id(0);                                           \n" \
"   int y = get_global_id(1);                                           \n" \
"   output[x][y] = image1[x][y] * IMAGE_SOURCE1_LIMPID + image2[x][y] * IMAGE_SOURCE2_LIMPID; \n" \
"}                                                                      \n" \
"\n";


  


int main( int argc, char* argv[] )  
{
    int len = 10;
    cl_mem d_a;
    cl_mem d_b;
    cl_mem d_c;
    cl_platform_id cpPlatform;        // OpenCL 平台
    cl_device_id device_id;           // device ID
    cl_context context,context1;               // context
    cl_command_queue queue,queue1;           // command queue
    cl_program program,program1;               // program
    cl_kernel kernel,kernel1;                 // kernel
    size_t bytes = len*sizeof(int);
    /*h_a = (int*)malloc(bytes);
    h_b = (int*)malloc(bytes);
    h_c = (int*)malloc(bytes);*/
    size_t globalSize, localSize;
    cl_int err;
    localSize = 2;
    globalSize = (size_t)ceil(len/(float)localSize)*localSize;
    





      
    float *image1, *image2;             // original data set given to device   
    float *results;                     // results returned from device   
    unsigned int correct;               // number of correct results returned   
  
    size_t global;                      // global domain size for our calculation   
    size_t local;                       // local domain size for our calculation   
    cl_mem input1, input2;              // device memory used for the input array   
    cl_mem output;                      // device memory used for the output array   
  
  
  
// Initialize the original data buffer and the result buffer   
    image1 = (float*)malloc(IMAGE_X_PIXELS * IMAGE_Y_PIXELS * sizeof(*image1));   
    image2 = (float*)malloc(IMAGE_X_PIXELS * IMAGE_Y_PIXELS * sizeof(*image2));   
    results = (float*)malloc(IMAGE_X_PIXELS * IMAGE_Y_PIXELS * sizeof(*results));   
    const unsigned int count = IMAGE_X_PIXELS * IMAGE_Y_PIXELS;   
  
// Automatically generate 2 images   
    for(int i = 0; i < count; i++) {   
        image1[i] = rand() / (float)RAND_MAX;   
        image2[i] = rand() / (float)RAND_MAX;   
    }   










    err = clGetPlatformIDs(1, &cpPlatform, NULL);
    err = clGetDeviceIDs(cpPlatform, CL_DEVICE_TYPE_CPU, 1, &device_id, NULL);
    context = clCreateContext(0, 1, &device_id, NULL, NULL, &err);
    context1 = clCreateContext(0, 1, &device_id, NULL, NULL, &err);

    queue = clCreateCommandQueue(context, device_id, 0, &err);
queue1 = clCreateCommandQueue(context1, device_id, 0, &err);
    program = clCreateProgramWithSource(context, 1, (const char **) & kernelSource, NULL, &err);
program1 = clCreateProgramWithSource(context1, 1, (const char **) & KernelSource1, NULL, &err);
    clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
    clBuildProgram(program1, 0, NULL, NULL, NULL, NULL);
    kernel = clCreateKernel(program, "vecMul", &err);
    kernel1 = clCreateKernel(program1, "Limpid", &err);

    d_a = clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, bytes, NULL, NULL);
    d_b = clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_ALLOC_HOST_PTR, bytes, NULL, NULL);
    d_c = clCreateBuffer(context, CL_MEM_WRITE_ONLY|CL_MEM_ALLOC_HOST_PTR, bytes, NULL, NULL);
    // （将向量信息写入设备缓冲）
    /*err = clEnqueueWriteBuffer(queue, d_a, CL_TRUE, 0,
        bytes, h_a, 0, NULL, NULL);
    err = clEnqueueWriteBuffer(queue, d_b, CL_TRUE, 0,
        bytes, h_b, 0, NULL, NULL);
    // （设置计算内核的参数）*/
    err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &d_a);
    err = clSetKernelArg(kernel, 1, sizeof(cl_mem), &d_b);
    err = clSetKernelArg(kernel, 2, sizeof(cl_mem), &d_c);
    err = clSetKernelArg(kernel, 3, sizeof(int), &len);



int *mappedBuffer_a = NULL;
int *mappedBuffer_b = NULL;
int *mappedBuffer_c = NULL;

mappedBuffer_a = (int *)clEnqueueMapBuffer(queue, d_a, CL_TRUE, CL_MAP_WRITE, 0, bytes, 0, NULL, NULL, NULL);
mappedBuffer_b = (int *)clEnqueueMapBuffer(queue, d_b, CL_TRUE, CL_MAP_WRITE, 0, bytes, 0, NULL, NULL, NULL);
mappedBuffer_c = (int *)clEnqueueMapBuffer(queue, d_c, CL_TRUE, CL_MAP_READ, 0, bytes, 0, NULL, NULL, NULL);
clFinish(queue);






// Create the input and output arrays in device memory for our calculation   
    //   
    input1 = clCreateBuffer(context1,  CL_MEM_READ_ONLY,  sizeof(float) * count, NULL, NULL);   
    input2 = clCreateBuffer(context1,  CL_MEM_READ_ONLY,  sizeof(float) * count, NULL, NULL);   
    output = clCreateBuffer(context1, CL_MEM_WRITE_ONLY, sizeof(float) * count, NULL, NULL);   
    if (!input1 || !input2 || !output)   
    {   
printf("Error: Failed to allocate device memory!\n");   
        exit(1);   
    }       
  
// Write our data set into the input array in device memory    
    //   
    err = clEnqueueWriteBuffer(queue1, input1, CL_TRUE, 0, sizeof(float) * count, image1, 0, NULL, NULL);   
    err |= clEnqueueWriteBuffer(queue1, input2, CL_TRUE, 0, sizeof(float) * count, image2, 0, NULL, NULL);   
    if (err != CL_SUCCESS)   
    {   
printf("Error: Failed to write to source array!\n");   
        exit(1);   
    }   
  
// Set the arguments to our compute kernel   
    //   
    err = 0;   
    err = clSetKernelArg(kernel1, 0, sizeof(cl_mem), &input1);   
    err |= clSetKernelArg(kernel1, 1, sizeof(cl_mem), &input2);   
    err |= clSetKernelArg(kernel1, 2, sizeof(cl_mem), &output);   
    if (err != CL_SUCCESS)   
    {   
printf("Error: Failed to set kernel arguments! %d\n", err);   
        exit(1);   
    }   
  
// Get the maximum work group size for executing the kernel on the device   
    //   
    err = clGetKernelWorkGroupInfo(kernel1, device_id, CL_KERNEL_WORK_GROUP_SIZE, sizeof(local), &local, NULL);   
    if (err != CL_SUCCESS)   
    {   
printf("Error: Failed to retrieve kernel work group info! %d\n", err);   
        exit(1);   
    }   
    else   
printf("The number of work items in a work group is: %lu\r\n", local);   
  
// Execute the kernel over the entire range of our 1d input data set   
// using the maximum number of work group items for this device   
    //   
    global = count; 






















for (int j=1;j<10;j++) {

     for( int i = 0; i < len; i++ )
    {
        mappedBuffer_a[i] = j+1;
	mappedBuffer_b[i] = j;
    }

    err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &globalSize, &localSize,0, NULL, NULL);
    clFinish(queue);
    /*clEnqueueReadBuffer(queue, d_c, CL_TRUE, 0,
        bytes, h_c, 0, NULL, NULL );
    float sum = 0;
    for(i=0; i<n; i++)
        printf("%d =====",h_c[i]);
    */

	 std::cout<<"+++c=(a*b)++++"<<std::endl;
         for( int i = 0; i < len; i++ )
    {
        std::cout<<mappedBuffer_c[i]<<"=("<< mappedBuffer_a[i]<<"*"<<mappedBuffer_b[i]<<") |  ";
    }
         std::cout<<"\n"<<"========"<<std::endl;


}





  err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &globalSize, &localSize,0, NULL, NULL);
   size_t globalSize1[2] = {IMAGE_X_PIXELS, IMAGE_Y_PIXELS};
   size_t  localSize1[2]={22, 12};
   err = clEnqueueNDRangeKernel(queue1, kernel1, 2, NULL, globalSize1, localSize1,0, NULL, NULL);
    if (err)   
    {   
printf("Error: Failed to execute kernel!\n");   
return EXIT_FAILURE;   
    }   
  
// Wait for the command commands to get serviced before reading back results   
    //   
    clFinish(queue1);   
  
// Read back the results from the device to verify the output   
    //   
    err = clEnqueueReadBuffer(queue1, output, CL_TRUE, 0, sizeof(float) * count, results, 0, NULL, NULL );     
    if (err != CL_SUCCESS)   
    {   
printf("Error: Failed to read output array! %d\n", err);   
        exit(1);   
    }   
  
// Validate our results   
    //   
    correct = 0;   
    for(int i = 0; i < count; i++)   
    {   
        if(results[i] == image1[i] * IMAGE_SOURCE1_LIMPID  + image2[i] * IMAGE_SOURCE2_LIMPID)   
            correct++;   
    }   
  
// Print a brief summary detailing the results   
    //   
printf("Computed '%d/%d' correct values!\n", correct, count);   









clEnqueueUnmapMemObject(queue, d_c, mappedBuffer_c, 0, NULL, NULL);
clEnqueueUnmapMemObject(queue, d_a, mappedBuffer_a, 0, NULL, NULL);
clEnqueueUnmapMemObject(queue, d_b, mappedBuffer_b, 0, NULL, NULL);
clFinish(queue);
    clReleaseMemObject(d_a);
    clReleaseMemObject(d_b);
    clReleaseMemObject(d_c);
    clReleaseProgram(program);
    clReleaseKernel(kernel);
    clReleaseCommandQueue(queue);
    clReleaseContext(context);
    /*free(h_a);
    free(h_b);
    free(h_c);
    return 0;*/

clReleaseMemObject(input1);   
clReleaseMemObject(input2);   
clReleaseMemObject(output);   
clReleaseProgram(program1);   
clReleaseKernel(kernel1);   
clReleaseCommandQueue(queue1);   
clReleaseContext(context1);   
  
free(image1);   
free(image2);   
free(results);   




}