下面对4x4矩阵进行转置,调用openCL,4X4矩阵采用二维数组进行存储,在程序设计上,让转置过程分10次转置完成,就是一次转一行,因此OpenCL的工作维数是二维。
openCL程序分为两个部份,一部份是内核代码,负责具体算法。另一部份是主程序负责初始化OpenCL和准备数据。主程序加载内核代码,并按照既定方法进行运算。
核函数代码如下
__kernel void matrix_transposition(__global int* a, __global int* b)
{
  int col = get_global_id(0);
  int row = get_global_id(1);
  b[col*4+row] = a[row*4+col];
}
主机端代码如下
int main()
{
  cl_device_id device;
  cl_platform_id platform_id = NULL;
  cl_context context;
  cl_command_queue cmdQueue;
  cl_mem bufferA,bufferB;
  cl_program program;
  cl_kernel kernel = NULL;
 
  size_t ocl_string_size;
  char *ocl_string;
  ocl_string = (char *)malloc(1024*1024);
  //我们使用的是二维向量
  //设定向量大小(维数)
  size_t globalWorkSize[2];
  globalWorkSize[0] = dim_x ;
  globalWorkSize[1] = dim_y;
  /*
  定义输入变量和输出变量,并设定初值
  */
  int buf_A[dim_x][dim_y];
  int buf_B[dim_x][dim_y];
  size_t datasize = sizeof(int) * dim_x * dim_y;
  int n=0;
  int m=0;
  for(n=0;n<dim_x;n++)
  {
  for(m=0;m<dim_y;m++)
  {
    buf_A[m][n] = m + n*dim_x;
  }
  }
  ocl_string_size = get_ocl_string("matrix_transposition.cl", ocl_string);
  //step 1:初始化OpenCL
  clGetPlatformIDs(1, &platform_id, NULL);
  clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_ACCELERATOR, 1, &device, NULL);
  //step 2:创建上下文
  context = clCreateContext(NULL,1,&device,NULL,NULL,NULL);
  //step 3:创建命令队列
  cmdQueue = clCreateCommandQueue(context,device,0,NULL);
  //step 4:创建数据缓冲区
  bufferA = clCreateBuffer(context,
         CL_MEM_READ_ONLY,
         datasize,NULL,NULL);
  bufferB = clCreateBuffer(context,
         CL_MEM_WRITE_ONLY,
         datasize,NULL,NULL);
  //step 5:将数据上传到缓冲区
  clEnqueueWriteBuffer(cmdQueue,
         bufferA,CL_FALSE,
         0,datasize,
         buf_A,0,
         NULL,NULL);
  //step 6:加载编译代码,创建内核调用函数
  program = clCreateProgramWithSource(context,1,
              (const char**)&ocl_string,
              &ocl_string_size,NULL);
  clBuildProgram(program,1,&device,NULL,NULL,NULL);
  kernel = clCreateKernel(program,"matrix_transposition",NULL);
  //step 7:设置参数,执行内核
  clSetKernelArg(kernel,0,sizeof(cl_mem),&bufferA);
  clSetKernelArg(kernel,1,sizeof(cl_mem),&bufferB);
  //<span style="color: #ff0000;"><strong>注意这里第三个参数已经改成2,表示二维数据。</strong></span>
  clEnqueueNDRangeKernel(cmdQueue,kernel,
         2,NULL,
         globalWorkSize,
         NULL,0,NULL,NULL);
  //step 8:取回计算结果
  clEnqueueReadBuffer(cmdQueue,bufferB,CL_TRUE,0,
        datasize,buf_B,0,NULL,NULL);
  //输出计算结果
  for(n=0;n<dim_x;n++)
  {
  for(m=0;m<dim_y;m++)
  {
    printf("%d ", buf_A[m][n]);
  }
  printf("\n");
  }
  for(n=0;n<dim_x;n++)
  {
  for(m=0;m<dim_y;m++)
  {
    printf("%d ", buf_B[m][n]);
  }
  printf("\n");
  }
  //释放所有调用和内存
  clReleaseKernel(kernel);
  clReleaseProgram(program);
  clReleaseCommandQueue(cmdQueue);
  clReleaseMemObject(bufferA);
  clReleaseMemObject(bufferB);
  clReleaseContext(context);
  return 0;
}
将上述程序编译,结果如下
0 1 2 3
4 5 6 7
8 9 10 11
12 13 14 25
矩阵转置后
0 4 8 12
1 5 9 13
2 6 10 14
3 7 11 15
该矩阵的行和列转换成功,即第 n 行变成第 n 列