OpenCL: OpenCL快速入门教程 |
您所在的位置：网站首页 › c语言硬件编程教程pdf › OpenCL: OpenCL快速入门教程 |
OpenCL: 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框架管理下的若干设备构成了这个平台，通过这个平台，应用程序可以与设备共享资源并在设备上执行kernel。平台通过cl_plantform来展现，可以使用下面的代码来初始化平台： // Returns the error code cl_int oclGetPlatformID (cl_platform_id *platforms) // Pointer to the platform objectDevice（设备）：通过cl_device来表现，使用下面的代码： // Returns the error code cl_int clGetDeviceIDs (cl_platform_id platform, cl_device_type device_type, // Bitfield identifying the type. For the GPU we use CL_DEVICE_TYPE_GPU cl_uint num_entries, // Number of devices, typically 1 cl_device_id *devices, // Pointer to the device object cl_uint *num_devices) // Puts here the number of devices matching the device_type
Context（上下文）：定义了整个OpenCL化境，包括OpenCL kernel、设备、内存管理、命令队列等。上下文使用cl_context来表现。使用以下代码初始化： // Returs the context cl_context clCreateContext (const cl_context_properties *properties, // Bitwise with the properties (see specification) cl_uint num_devices, // Number of devices const cl_device_id *devices, // Pointer to the devices object void (*pfn_notify)(const char *errinfo, const void *private_info, size_t cb, void *user_data), // (don't worry about this) void *user_data, // (don't worry about this) cl_int *errcode_ret) // error code resultCommand-Queue（指令队列）：就像它的名字一样，他是一个存储需要在设备上执行的OpenCL指令的队列。“指令队列建立在一个上下文中的指定设备上。多个指令队列允许应用程序在不需要同步的情况下执行多条无关联的指令。” cl_command_queue clCreateCommandQueue (cl_context context, cl_device_id device, cl_command_queue_properties properties, // Bitwise with the properties cl_int *errcode_ret) // error code result下面的例子展示了这些元素的使用方法： cl_int error = 0; // Used to handle error codes cl_platform_id platform; cl_context context; cl_command_queue queue; cl_device_id device; // Platform error = oclGetPlatformID(&platform); if (error != CL_SUCCESS) { cout |
【本文地址】