OpenCL: OpenCL快速入门教程

分享一下我老师大神的人工智能教程！零基础，通俗易懂！http://blog.csdn.net/jiangjunshow

也欢迎大家转载本篇文章。分享知识，造福人民，实现我们中华民族伟大复兴！

               

OpenCL快速入门教程

原文地址：http://opencl.codeplex.com/wikipage?title=OpenCL%20Tutorials%20-%201

翻译日期：2012年6月4日星期一

 

 这是第一篇真正的OpenCL教程。这篇文章不会从GPU结构的技术概念和性能指标入手。我们将会从OpenCL的基础API开始，使用一个小的kernel作为例子来讲解基本的计算管理。

 首先我们需要明白的是，OpenCL程序是分成两部分的：一部分是在设备上执行的（对于我们，是GPU），另一部分是在主机上运行的（对于我们，是CPU）。在设备上执行的程序或许是你比较关注的。它是OpenCL产生神奇力量的地方。为了能在设备上执行代码，程序员需要写一个特殊的函数（kernel函数）。这个函数需要使用OpenCL语言编写。OpenCL语言采用了C语言的一部分加上一些约束、关键字和数据类型。在主机上运行的程序提供了API，所以i可以管理你在设备上运行的程序。主机程序可以用C或者C++编写，它控制OpenCL的环境（上下文，指令队列…）。

设备（Device）

 我们来简单的说一下设备。设备，像上文介绍的一样，OpenCL编程最给力的地方。

 我们必须了解一些基本概念：

 Kernel：你可以把它想像成一个可以在设备上执行的函数。当然也会有其他可以在设备上执行的函数，但是他们之间是有一些区别的。Kernel是设备程序执行的入口点。换言之，Kernel是唯一可以从主机上调用执行的函数。

 现在的问题是：我们如何来编写一个Kernel？在Kernel中如何表达并行性？它的执行模型是怎样的？解决这些问题，我们需要引入下面的概念：

     SIMT：单指令多线程（SINGLE INSTRUCTION MULTI THREAD）的简写。就像这名字一样，相同的代码在不同线程中并行执行，每个线程使用不同的数据来执行同一段代码。

     Work-item（工作项）：Work-item与CUDA Threads是一样的，是最小的执行单元。每次一个Kernel开始执行，很多（程序员定义数量）的Work-item就开始运行，每个都执行同样的代码。每个work-item有一个ID，这个ID在kernel中是可以访问的，每个运行在work-item上的kernel通过这个ID来找出work-item需要处理的数据。

     Work-group（工作组）：work-group的存在是为了允许work-item之间的通信和协作。它反映出work-item的组织形式（work-group是以N维网格形式组织的，N=1，2或3）。

 Work-group等价于CUDA thread blocks。像work-items一样，work-groups也有一个kernel可以读取的唯一的ID。

 ND-Range:ND-Range是下一个组织级别，定义了work-group的组织形式（ND-Rang以N维网格形式组织的，N=1，2或3）；

 这是ND-Range组织形式的例子

Kernel

 现在该写我们的第一个kernel了。我们写一个小的kernel将两个向量相加。这个kernel需要四个参数：两个要相加的向量，一个存储结果的向量，和向量个数。如果你写一个程序在cpu上解决这个问题，将会是下面这个样子：

void vector_add_cpu (const float* src_a,               const float* src_b,               float*  res,               const int num){   for (int i = 0; i < num; i++)      res[i] = src_a[i] + src_b[i];}

 在GPU上，逻辑就会有一些不同。我们使每个线程计算一个元素的方法来代替cpu程序中的循环计算。每个线程的index与要计算的向量的index相同。我们来看一下代码实现：

__kernel void vector_add_gpu (__global const float* src_a,                     __global const float* src_b,                     __global float* res,           const int num){   /* get_global_id(0) 返回正在执行的这个线程的ID。   许多线程会在同一时间开始执行同一个kernel，   每个线程都会收到一个不同的ID，所以必然会执行一个不同的计算。*/   const int idx = get_global_id(0);   /* 每个work-item都会检查自己的id是否在向量数组的区间内。   如果在，work-item就会执行相应的计算。*/   if (idx < num)      res[idx] = src_a[idx] + src_b[idx];}

有一些需要注意的地方：

1. Kernel关键字定义了一个函数是kernel函数。Kernel函数必须返回void。

2. Global关键字位于参数前面。它定义了参数内存的存放位置。

 另外，所有kernel都必须写在“.cl”文件中，“.cl”文件必须只包含OpenCL代码。

主机（Host）

 我们的kernel已经写好了，现在我们来写host程序。

建立基本OpenCL运行环境

 有一些东西我们必须要弄清楚：

Plantform（平台）：主机加上OpenCL框架管理下的若干设备构成了这个平台，通过这个平台，应用程序可以与设备共享资