<span style="font-size:10px;">_kernel void vecadd(_global int *C, _global int *A, _global int *C)
{int tid = get_global_id(0);    //OpenCL intrisic functionC[tid] = A[tid] + B[tid];
}</span>

1. OpenCL源代码以字符串形式存储。如果代码存放在文件中，则必须读取到内存中，在内存中以字符串数组的形式存放。
2. 源代码通过调用clCreateProgramWithSource()转化为cl_program对象。
3. 调用clBuildProgram()在支持OpenCL的设备上编译程序对象。

向量加代码演示

#include<stdio.h>
#include<stdlib.h>
#include<CL/cl.h>const char *programSource =
"_kernel void vecadd(_global int *A, _global int *B, _global int *C) \n"
"{                                                                   \n"
"  int idx = get_global_id(0);                                  \n"
"  C[idx] = A[idx] + B[idx];                                    \n"
"}                                                                   \n"    int main()
{int *A = NULL;int *B = NULL;int *C = NULL;const int elements = 2048;size_t datasize = sizeof(int) * elements;A = (int*)malloc(datasize);B = (int*)malloc(datasize);C = (int*)malloc(datasize);int i;for(i=0; i<elements; i++){A[i] = i;B[i] = i;}//Check return valuecl_int status;//Retrieve the number of platforms cl_uint numPlatforms = 0;status = clGetPlatformIDs(0, NULL, &numPlatforms);//Allocate space for each platformcl_plarform_id *platforms = NULL;platforms =  (cl_platform_id*)malloc(numPlatforms * sizeof(cl_platform_id));//Fill in the platformsstatus = clGetPlatformIDs(numPlatforms, platforms, NULL);//Retrieve the number of devicescl_unit numDevicdes = 0;status = clGetDeviceIDs(platform[0], CL_DEVICE_TYPE_ALL, 0, NULL, &numDevices);//Allocate space for devicescl_device_id *devices;devices = (cl_device_id*)malloc(numDevices * sizeof(cl_device_id));//Fill in the devicesstatus = clGetDeviceIDs(platform[0], CL_DEVICE_TYPE_ALL, numDevices, deivces, NULL);//Create a context and  associate it with devicescl_context context;context = clCreateContext(NULL, numDevices, devices, NULL, NULL, &status);//Create a command queue and associate it with the devicecl_command_queue cmdQueue;cmdQueue = clCreateCommandQueue(context, devices[0], 0, &status);//Create buffer objectcl_mem bufA = clCreateBuffer(context, CL_MEM_READ_ONLY, datasize, NULL, &status);cl_mem bufB = clCreateBuffer(context, CL_MEM_READ_ONLY, datasize, NULL, &status);cl_mem bufC = clCreateBuffer(context, CL_MEM_READ_ONLY, datasize, NULL, &status);//Write to the device bufferstatus = clEnqueueWriteBuffer(cmdQueue, bufA, CL_FALSE, 0, datasize, A, 0, NULL, NULL);status = clEnqueueWriteBuffer(cmdQueue, bufB, CL_FALSE, 0, datasize, B, 0, NULL, NULL);status = clEnqueueWriteBuffer(cmdQueue, bufC, CL_FALSE, 0, datasize, C, 0, NULL, NULL);//Create program from source codecl_program program = clCreateProgramWithSource(context, 1, (const char**)&programSource, NULL, &status);//Compile the programstatus = clBuildProgram(program, "vacadd", &status);//Associate buffers with the kernelstatus = clSetKernelArg(kernel, 0, sizeof(cl_mem), &bufA);status = clSetKernelArg(kernel, 1, sizeof(cl_mem), &bufB);status = clSetKernelArg(kernel, 2, sizeof(cl_mem), &bufC);//Define index spacesize_t globalWorkSize[1];globalWorkSize[0] = elements;//execute the kernelstatus = clEnqueueNDRangeKernel(cmdQueue, kernle, 1, NULL, globalWorkSize, NULL, 0, NULL, NULL);//Read the device output buffer to the host output array clEnqueueReadBuffer(cmdQueue, bufC, CL_TRUE, 0, datasize, C, 0, NULL, NULL);//Verify the outputint result = 1;for(i=0;i<elements;i++){if(C[i] != i+i){result = 0;break;}}if(result){printf("correct!\n");}else{printf("incorrect!\n");}//Release resourcesclReleaseKernel(kernel);clReleaseProgram(program);clReleaseCommandQueue(cmdQueue);clReleaseMemObject(bufA);clReleaseMemObject(bufB);clReleaseMemObject(bufC);clReleaseContext(context);//Free host resourcesfree(A);free(B);free(C);free(platform);free(devices);
}