摘要：本篇教程中，我们学习一下如何用opencl有效实现数组求和，也就是通常所说的reduction问题。 在程序中，我们设置workgroup size为256，kernel的输入、输出缓冲参数都用uint4的格式，这样我们原始求和的数组大小为256*4的倍数，数据类型为uint。我们设定每个workgroup处理处理512个uint4，即2048个uint 为了简便... 阅读全文

摘要：本章学习一下在opencl中如何实现矩阵的转置，主要的技巧还是利用好local memory，防止bank conflit以及使得全局memory的读写尽量是合并(coalensing)读写。 我们的矩阵是一副二维灰度图像256*256，矩阵的转置也就是图像的转置。每个thread处理16(4*4)个pixel(uchar)，workgroup的size是(16,16)。 ... 阅读全文

摘要：在OpenCL编程中，特别是基于GPU的opencl的编程，提高程序性能最主要的方法就是想法提高memory的利用率，它主要包括两方面的优化：一方面是提高global memory的合并读写效率，另一方面就是减少local memory的bank conflict。下面我们分析一下教程7中的代码，看看它的memory利用率如何？ 首先我们用amd的opencl profil... 阅读全文

摘要：现在我们利用上一篇教程的方法，来统计一副RGBA图像中有多少个像素点（该像素点满足R, G, B, A任意分量>=5)。我考虑的方法是建立256 bin的直方图，对于一个像素，求max(R, G,B,A),用该值决定该像素点进入那个bin，这样求出直方图后，width*height - hostBi 阅读全文

摘要：histogram翻译成中文就是直方图，在计算机图像处理和视觉技术中，通常用histogram来进行图像匹配，从而完成track，比如meanshift跟踪算法中，经常要用到图像的直方图。 灰度图的histogram计算，首先要选择bin(中文可以称作槽)的数量，对于灰度图，像素的范围通常是[0-255],所以bin的数目就是256，然后我们循环整幅图像，统计出每种像素值出现... 阅读全文

摘要：在本教程中，我们学习用opencl进行简单的图像处理，对一个图片进行旋转。图片读入、保存等工作，我们使用开源的FreeImage，下载地址：http://freeimage.sourceforge.net/ 首先我们建立一个gFreeImage类，用来装入图像，该类主要调用FreeImage的函数，首先会初始化FreeImage库，然后根据文件名猜测图像文件格式，最终lo... 阅读全文

摘要：在本教程中，我们使用二维NDRange来设置workgroup，这样在opencl中，workitme的组织形式是二维的，Kernel中 的代码也要做相应的改变，我们先看一下clEnqueueNDRangeKernel函数的变化。首先我们指定了workgroup size为localx*localy,通常这个值为64的倍数，但最好不要超过256。 //执行kernel,Range用2... 阅读全文

摘要：本教程中，我们使用上一篇教程中产生的二进制kernel文件vecadd.bin作为输入来创建程序对象，程序代码如下： //kernel文件为vecadd.bin gclFile kernelFile; if(!kernelFile.readBinaryFromFile("vecadd.bin")) { printf("Failed... 阅读全文

摘要：在教程二中，我们通过函数convertToString，把kernel源文件读到一个string串中，然后用函数clCreateProgramWithSource装入程序对象，再调用函数clBuildProgram编译程序对象。其实我们也可以直接调用二进制kernel文件，这样，当不想把kernel文件给别人看的时候，起到一定的保密作用。在本教程中，我们会把读入的源文件存储一个二进制文件... 阅读全文

摘要：现在，我们开始写一个简单的OpenCL程序，计算两个数组相加的和，放到另一个数组中去。程序用cpu和gpu分别计算，最后验证它们是否相等。OpenCL程序的流程大致如下： 下面是source code中的主要代码： int main(int argc, char* argv[]) { //在host内存中创建三个缓冲区 ... 阅读全文

摘要：从今天开始学习OpenCL，下面先介绍OpenCL APP(Accelerated Parallel processing)的安装。 下载地址：http://developer.amd.com/tools-and-sdks/opencl-zone/amd-accelerated-parallel-processing-app-sdk/ 根据你的OS，选择相应的SDK版本。 安装注... 阅读全文

摘要：在本节，我们主要介绍OpenCL中buffer的使用，同时提供了两个个完整的例子，一个是图像的旋转，一个是矩阵乘法（非常简单，没有分块优化）。 1、创建OpenCL设备缓冲(buffer) OpenCL设备使用的数据都存放在设备的buffer中[其实就是device memory中]。我们用下面的代码创建buffer对象： cl_mem bufferobj = clCreat... 阅读全文

摘要：1、OpenCL扩展 OpenCL扩展是指device支持某种特性，但这中特性并不是OpenCL标准的一部分。通过扩展，厂商可以给device增加一些新的功能，而不用考虑兼容性问题。现在各个厂商在OpenCL的实现中或多或少的使用了自己的扩展。 扩展的类型分为三种： Khronos OpenCL工作组批准的扩展，这种要经过一致性测试，可能会被增加到新版本的Open... 阅读全文

摘要：本节主要介绍NBody算法的OpenCL性能优化。 1、NBody NBody系统主要用来通过粒子之间的物理作用力来模拟星系系统。每个粒子表示一个星星，多个粒子之间的相互作用，就呈现出星系的效果。 上图为一个粒子模拟星系的图片：Source: THE GALAXY-CLUSTER-SUPERCLUSTER CONNECTION，http://www.casca... 阅读全文

摘要：性能优化 1、线程映射 所谓线程映射是指某个线程访问哪一部分数据，其实就是线程id和访问数据之间的对应关系。 合适的线程映射可以充分利用硬件特性，从而提高程序的性能，反之，则会降低性能。 请参考Static Memory Access Pattern Analysis on a Massively Parallel GPU这篇paper，文中讲述线程如何在算法中充分利用线程映射... 阅读全文

摘要：GPU线程及调度 本节主要讲述OpenCL中的Workgroup如何在硬件设备中被调度执行。同时也会讲一下同一个Workgroup中的workitem，如果它们执行的指令发生diverage（就是执行指令不一致）对性能的影响。学习OpenCL并行编程，不仅仅是对OpenCL Spec本身了解，更重要的是了解OpenCL硬件设备的特性，现阶段来说，主要是了解GPU的的架构特性，这样才能针... 阅读全文

摘要：本节主要讲述GPU的memory架构。优化基于GPU device的kernel程序时，我们需要了解很多GPU的memory知识，比如内存合并，bank conflit（冲突）等等，这样才能针对具体算法做一些优化工作。 1、GPU总线寻址介绍 假定X是一个指向整数（32位整数）数组的指针，数组的首地址为0x00001232。一个线程要访问元素X[0], int t... 阅读全文

摘要：6、Nvdia GPU Femi架构 GTX480-Compute 2.0 capability： 有15个core或者说SM（Streaming Multiprocessors ）。 每个SM,一般有32 cuda处理器。 共480个cuda处理器。 带ECC的global memory 每个SM内的线程按32个单位调度执行，称作warp。每... 阅读全文

摘要：GPU架构 内容包括： 1.OpenCLspec和多核硬件的对应关系 AMD GPU架构 Nvdia GPU架构 Cell Broadband Engine 2.一些关于OpenCL的特殊主题 OpenCL编译系统 Installable client driver 首先我们可能有疑问，既然OpenCL具有平台无关性，我们为什么还要去研... 阅读全文

摘要：OpenCL内存模型 OpenCL的内存模型定义了各种各样内存类型，各种内存模型之间有层级关系。各种内存之间的数据传输必须是显式进行的，比如从host memory到device memory，从global memory到local memory等等。 WorkGroup被映射到硬件的CU上执行（在AMD 5xxx系列显卡上，CU就是simd，一个simd中有16个p... 阅读全文