opencl.cpp

#include <stdio.h>
#include <stdlib.h>

#ifdef __APPLE__
  #include <OpenCL/opencl.h>
#else
  #include <CL/cl.h>
#endif

#define MAX_SOURCE_SIZE (0x100000)

int main ( void )
{
  int i;
  const int LIST_SIZE = 1024;
  int* A = ( int* ) malloc ( sizeof ( int ) * LIST_SIZE );
  int* B = ( int* ) malloc ( sizeof ( int ) * LIST_SIZE );

  for ( i = 0; i < LIST_SIZE; i++ )
  {
    A[i] = i;
    B[i] = LIST_SIZE - i;
  }

  //载入内核源码到source_str
  FILE* fp;
  char* source_str;
  size_t source_size;

  fp = fopen ( "vector_add_kernel.cl", "r" );

  if ( !fp )
  {
    fprintf ( stderr, "Failed to load kernel\n" );
    exit ( 1 );
  }

  source_str = ( char* ) malloc ( MAX_SOURCE_SIZE );
  source_size = fread ( source_str, 1, MAX_SOURCE_SIZE, fp );
  fclose ( fp );

  //获得平台和设备信息
  cl_platform_id platform_id = NULL;
  cl_device_id device_id = NULL;
  cl_uint ret_num_devices;
  cl_uint ret_num_platforms;
  cl_int ret = clGetPlatformIDs ( 1, &platform_id, &ret_num_platforms );
  ret = clGetDeviceIDs ( platform_id, CL_DEVICE_TYPE_DEFAULT, 1, &device_id, &ret_num_devices );

  //创建OpenCL上下文
  cl_context context = clCreateContext ( NULL, 1, &device_id, NULL, NULL, &ret );
  //创建命令队列
  cl_command_queue command_queue = clCreateCommandQueue ( context, device_id, 0, &ret );
  //创建内存缓冲对象，在设备上为每个向量
  cl_mem a_mem_obj = clCreateBuffer ( context, CL_MEM_READ_ONLY,
                                      LIST_SIZE * sizeof ( int ), NULL, &ret );
  cl_mem b_mem_obj = clCreateBuffer ( context, CL_MEM_READ_ONLY,
                                      LIST_SIZE * sizeof ( int ), NULL, &ret );
  cl_mem c_mem_obj = clCreateBuffer ( context, CL_MEM_WRITE_ONLY,
                                      LIST_SIZE * sizeof ( int ), NULL, &ret );

  //拷贝数据A和B到对应的内存缓冲
  ret = clEnqueueWriteBuffer ( command_queue, a_mem_obj, CL_TRUE, 0,
                               LIST_SIZE * sizeof ( int ), A, 0, NULL, NULL );
  ret = clEnqueueWriteBuffer ( command_queue, a_mem_obj, CL_TRUE, 0,
                               LIST_SIZE * sizeof ( int ), B, 0, NULL, NULL );

  //创建程序
  cl_program program = clCreateProgramWithSource ( context, 1,
                       ( const char** ) &source_str, ( const size_t* ) &source_size, &ret );
  //构建程序
  ret = clBuildProgram ( program, 1, &device_id, NULL, NULL, NULL );
  //创建OpenCL内核
  cl_kernel kernel = clCreateKernel ( program, "vector_add", &ret );

  //设置内核参数
  ret = clSetKernelArg ( kernel, 0, sizeof ( cl_mem ), ( void* ) &a_mem_obj );
  ret = clSetKernelArg ( kernel, 1, sizeof ( cl_mem ), ( void* ) &b_mem_obj );
  ret = clSetKernelArg ( kernel, 2, sizeof ( cl_mem ), ( void* ) &c_mem_obj );

  //执行内核
  size_t global_item_size = LIST_SIZE; //处理整个列表
  size_t local_item_size = 64;        //分割为64个组
  ret = clEnqueueNDRangeKernel ( command_queue, kernel, 1, NULL,
                                 &global_item_size, &local_item_size, 0, NULL, NULL );

  //读取内存缓冲C到本地变量C
  int* C = ( int* ) malloc ( sizeof ( int ) * LIST_SIZE );
  ret = clEnqueueReadBuffer ( command_queue, c_mem_obj, CL_TRUE, 0,
                              LIST_SIZE * sizeof ( int ), C, 0, NULL, NULL );

  //显示结果
  for ( i = 0; i < LIST_SIZE; i++ )
  {
    printf ( "%d + %d = %d\n", A[i], B[i], C[i] );
  }

  //清理资源
  ret = clFlush ( command_queue );
  ret = clFinish ( command_queue );
  ret = clReleaseKernel ( kernel );
  ret = clReleaseProgram ( program );
  ret = clReleaseMemObject ( a_mem_obj );
  ret = clReleaseMemObject ( b_mem_obj );
  ret = clReleaseMemObject ( c_mem_obj );
  ret = clReleaseCommandQueue ( command_queue );
  ret = clReleaseContext ( context );
  free ( A );
  free ( B );
  free ( C );
  return 0;
}