简介
第一篇文章，OpenCL：连接平行世界的桥梁，是对OpenCL主题的简要介绍。它解决了opencl（也被称为内核，虽然不是很精确）中的程序与MQL5中的外部（主机）程序之间的交互的基本问题。通过计算pi=3.14159265，证明了某些语言（如使用矢量数据类型）的性能。
在某些情况下，程序的性能有相当大的优化空间。但是，以上所有的优化都是幼稚的，因为它们没有考虑执行所有计算所需的硬件规格。在大多数情况下，对这些规范的了解将清楚地使我们能够实现显著超过CPU性能的加速。
为了讨论这种优化，作者必须求助于一个不再是原始的例子，可能是OpenCL文献中研究得最充分的例子。即两个大矩阵的乘法。
我们将讨论opencl存储模型及其在实际硬件体系结构上的独特实现。
1。现代计算设备
的存储水平
1.1。OpenCL存储型号
一般来说，根据计算机平台的不同，存储系统的差异很大。例如，所有现代CPU都支持自动数据缓存，而GPU并不总是支持自动数据缓存。
为了保证代码的可移植性，OpenCL采用了一种抽象的存储模型，程序员和提供者可以在实际硬件上实现该模型。理论上，opencl中定义的存储如下图所示：
图1。OpenCL存储型号
一旦数据从主机传输到设备，它就存储在全局设备内存中。以相反方向传输的任何数据也将存储在全局内存中（但这次存储在全局主机内存中）。关键字_global（双下划线！）表示与特定指针关联的数据已存储在全局内存中的修饰符：
__kernel void foo( __global float *A ) { /// kernel code }
设备中的所有计算单元都可以访问全局内存，就像主机系统中的RAM一样。
与名称不同，常量内存不一定存储只读数据。这种存储类型是为所有工作单元都可以同时访问每个元素的数据而设计的。具有常量值的变量也包括在这一类别中。OpenCL模型中的常量内存是全局内存的一部分，因此传输到全局内存的存储对象可以指定为_常量。
本地存储器是指专用于每个设备的高速临时存储器。在硬件上，它通常以片上存储器的形式存在，但对OpenCL没有完全相同的具体要求。
访问这样的存储会导致更低的延迟，从而比全局存储更大的存储带宽。我们将尝试充分利用它的低延迟来优化内核性能。
根据opencl规范，本地内存中的变量可以在内核头
中声明。
__kernel void foo( __local float *sharedData ) { }
亦可在其主体内声明：
__kernel void foo( __global float *A ) { __local float sharedData[ 64 ]; }
注意：动态数组不能在内核主体中声明；您始终都要指定其大小。
在下面的两个大型矩阵相乘的内核优化过程中，您可以看到如何处理本地数据以及作者在MetaTrader 5中体验到的实现特性。
默认情况下，不包含指针的局部变量和内核参数是私有的（例如，在没有本地修饰符的情况下指定）。实际上，这些变量通常位于寄存器中。相反，私有阵列和任何溢出寄存器通常位于片外存储器中，即高延迟存储器。我们引用了一些关于维基百科的信息：
在许多编程语言中，程序员对任意配置许多变量有误解。但是在编译时，编译器必须决定如何在一个寄存器组较小且有限的系统中分配这些变量。并非所有变量都同时使用（或“活动”），因此某些寄存器可以分配多个变量。但是，如果同时使用两个变量，则不可能在不破坏它们的值的情况下将它们分配到同一个寄存器。
不能分配给同一个寄存器的变量必须保存在RAM（随机存取存储器）中，并在需要读或写时加载，这一过程称为溢出。RAM访问速度明显低于寄存器访问速度，这将降低编译器的运行速度，因此经过优化的编译器将尽可能多地将变量放入寄存器。当硬件寄存器的数目小于理想数目时，使用寄存器压力这个术语。高压情况通常意味着更多的溢流和过载。
寄存器压力是 GPU 编程要面对的现实，由于要在有限的芯片区域上布设大量的核心，所以不可能再放置许多的寄存器。
我们称之为
的OpenCL内存模型与现代GPU的存储结构非常相似。下图显示了OpenCL内存模型与GPU AMD Radeon HD 6970内存模型之间的关联。
图2。Radeon HD 6970存储结构与抽象OpenCL内存模型之间的关联
我们继续研究与特定GPU内存实现相关的更详细的问题。
1.2。现代独立GPU中的存储器
1.2.1合并内存请求
这些信息对于优化内核的性能也很重要，因为其主要目标是实现高存储带宽。
请看下图以更好地了解内存寻址过程：
图3。全局设备存储器中的地址数据图
假设指向整型变量int数组的指针是地址=0x0001232。每个int占用4个字节的内存。假设线程的地址数据（执行内核代码的工作单元的软件模拟）位于[0]：
int tmp = X[ 0 ];
我们假设内存总线宽度为32字节（256位）。这种总线宽度在强大的GPU（如Radeon HD 5870）中很常见。其他GPU可能具有不同的数据总线宽度，例如一些具有384位甚至512位的Nvidia型号。
存储总线的寻址应与其结构相对应，即其宽度是最重要的。换句话说，内存中的数据存储在每个32字节（256位）块中。无论我们如何在0x0001220到0x000123F的范围内寻址（都在32个精确字节的范围内，您可以自己查看），我们都会得到地址0x0001220作为读取的起始地址。
访问地址0x0001232返回0x0001220到0x000123F范围内的所有数据，即八个整数。因此，只有四个字节（整数）的有用数据，而其余28个字节（七个整数）的无用数据：
图4。用于从内存获取所需数据的架构图
我们在之前指定地址需要的号码-0x0001232-包含在示意图中。
为了最大限度地利用这个总线，GPU试图将通过不同线程访问的内存合并到一个存储请求中；存储访问越少越好。这背后的原因是访问全局设备内存占用了我们的时间，从而极大地影响了程序的速度。看看内核代码的下一行：
int tmp = X[ get_global_id( 0 ) ];
假设我们的数组来自前面给出的示例。前16个线程（内核）然后访问从0x0001232到0x0001272的地址（在此范围内有16个整数或64个字节）。如果每个请求都是通过内核独立发送的，不需要预先合并到单个内核中，那么16个请求中的每个请求将包含4个有用数据字节和28个无用数据字节，因此总共使用64个字节和448个未使用的字节。
此计算基于这样一个事实：对同一个32字节存储块中的地址的每次访问都返回绝对一致的数据。这是钥匙。将多个请求合并为一个一致的请求更准确，从而减少无用请求的数量。此操作在下文中称为合并，而合并请求则称为一致性。
图5。只有三个存储请求
需要获得所需数据
上图中的每个单元格都是4个字节。在这种情况下，三个请求就足够了。如果按地址将数组的开头与每个32字节块的开头对齐，则即使两个请求也足够。
在AMD GPU 64中，线程是波前的一部分，因此应按照SIMD执行中的相同指令执行。由GET_Global_ID（0）排列的16个线程（仅为波前的四分之一）被合并成一个一致的请求，以实现总线的有效使用。
下图显示了一致性请求和非一致性请求（即“自发”请求）所需的存储带宽之间的比较。此处使用的是Radeon HD 5870，但Nvidia图形卡显示的结果类似。
图6。相干和非相干请求
所需内存带宽的比较分析
很明显，相干存储请求允许存储带宽增加一个数量级。
1.2.2.仓库
内存由实际存储数据的库组成。在现代GPU中，它通常是32位（4字节）字。串行数据存储在相邻的存储库中。访问串行元素的线程组没有任何库冲突。
库冲突的最大负面影响出现在本地GPU内存中。因此，最好通过在不同的存储库中定位相邻的线程来访问本地数据。
在AMD硬件上，库冲突的波前不会停止，直到所有本地内存操作完成。这会导致序列化，其中要并行执行的代码块是串行执行的。它对内核的性能有非常负面的影响。
图7。无库冲突的存储访问方案图
上图显示了没有库冲突的存储访问，因为所有线程都访问不同的数据。
让我们再看一次与库冲突的存储访问：
图8。存储访问
与库冲突
但是也有例外：如果所有的访问都是同一个地址，库可以执行广播以避免延迟：
图9。所有线程访问同一地址
访问全局内存时会发生类似的事件，但这种冲突的影响很小。
1.2.3. GPU 内存：总结
- GPU内存不同于CPU内存。OpenCL性能优化的主要目标是确保最大带宽，而不是像CPU那样缩短延迟。
- 存储访问的性质对总线利用效率有很大的影响。低总线利用率意味着低运行速度。
- 为了提高代码性能，存储访问是最好的一致性。此外，最好避免图书馆冲突。
- 硬件规格（总线宽度、存储库数量和可以合并到单个一致访问中的线程数量）可以在供应商提供的相关文档中找到。
下面列出的Rideon 5xx系列中的一些图形卡规格示例是
0
图10。中高端Radeon HD 58xx图形卡技术规范
现在，让我们继续编程。
2。大块矩阵乘法：从串行CPU码到并行GPU码
2.1。MQL5代码
与前面的文章opencl:bridges连接平行世界不同，手头的标准任务是将两个矩阵相乘。选择它的原因是可以从不同的来源找到关于这个主题的大量信息。它们中的大多数以某种方式提供或多或少协调的解决方案。这就是我们要做的，一步一步地阐明模型结构的意义，同时记住我们要面对的是实际的硬件。
下面是线性代数领域的一个著名的矩阵乘法公式，它是专门为计算机操作而修改的。第一个索引是矩阵的行号，第二个索引是列号。每个输出矩阵元素通过依次将第一个和第二个矩阵中每个元素的连续积添加到累积和中来计算。最后，这个累积和是计算出的输出矩阵元素：
1
图11。矩阵乘法公式
数字如下：
2
图12。矩阵乘法算法图解（以输出矩阵元素的计算为例）
很容易看出，只要两个矩阵的维数都等于n，加法和乘法的次数就可以用函数o（n^3）来估计。要计算每个输出矩阵元素，需要获得第一个矩阵中一行和第二个矩阵中一列的无向积。加和乘大约需要2*n次。所需的估计值是通过乘以矩阵元素n^2得到的。因此，粗略的代码运行时间在很大程度上依赖于N.
的强大功能。
为了方便起见，矩阵的行和列被设置为2000；它们可以是任意数字，但不能太大。
MQL5中的代码并不十分复杂：
//+------------------------------------------------------------------+ //| matr_mul_2dim.mq5 | //+------------------------------------------------------------------+ #define ROWS1 1000 // rows in the first matrix #define COLSROWS 1000 // columns in the first matrix = rows in the second matrix #define COLS2 1000 // columns in the second matrix float first[ ROWS1 ][ COLSROWS ]; // first matrix float second[ COLSROWS ][ COLS2 ]; // second matrix float third[ ROWS1 ][ COLS2 ]; // product //+------------------------------------------------------------------+ //| Script program start function | //+------------------------------------------------------------------+ void OnStart() { MathSrand(GetTickCount()); Print("======================================="); Print("ROWS1 = "+i2s(ROWS1)+"; COLSROWS = "+i2s(COLSROWS)+"; COLS2 = "+i2s(COLS2)); genMatrices(); ArrayInitialize(third,0.0f); //--- execution on the CPU uint st1=GetTickCount(); mul(); double time1=(double)(GetTickCount()-st1)/1000.; Print("CPU: time = "+DoubleToString(time1,3)+" s."); return; } //+------------------------------------------------------------------+ //| i2s | //+------------------------------------------------------------------+ string i2s(int arg) { return IntegerToString(arg); } //+------------------------------------------------------------------+ //| genMatrices | //| generate initial matrices; this generation is not reflected | //| in the final runtime calculation | //+------------------------------------------------------------------+ void genMatrices() { for(int r=0; r<ROWS1; r++) for(int c=0; c<COLSROWS; c++) first[r][c]=genVal(); for(int r=0; r<COLSROWS; r++) for(int c=0; c<COLS2; c++) second[r][c]=genVal(); return; } //+------------------------------------------------------------------+ //| genVal | //| generate one value of the matrix element: | //| uniformly distributed value lying in the range [-0.5; 0.5] | //+------------------------------------------------------------------+ float genVal() { return(float)(( MathRand()-16383.5)/32767.); } //+------------------------------------------------------------------+ //| mul | //| Main matrix multiplication function | //+------------------------------------------------------------------+ void mul() { // r-cr-c: 10.530 s for(int r=0; r<ROWS1; r++) for(int cr=0; cr<COLSROWS; cr++) for(int c=0; c<COLS2; c++) third[r][c]+=first[r][cr]*second[cr][c]; return; }
清单1。主机上的初始顺序程序
不同参数的性能结果：
2012.05.19 09:39:11 matr_mul_2dim (EURUSD,H1) CPU: time = 10.530 s. 2012.05.19 09:39:00 matr_mul_2dim (EURUSD,H1) ROWS1 = 1000; COLSROWS = 1000; COLS2 = 1000 2012.05.19 09:39:00 matr_mul_2dim (EURUSD,H1) ======================================= 2012.05.19 09:41:04 matr_mul_2dim (EURUSD,H1) CPU: time = 83.663 s. 2012.05.19 09:39:40 matr_mul_2dim (EURUSD,H1) ROWS1 = 2000; COLSROWS = 2000; COLS2 = 2000 2012.05.19 09:39:40 matr_mul_2dim (EURUSD,H1) =======================================
从
可以看出，我们对线性矩阵大小依赖性的运行时估计似乎是正确的：将所有矩阵的维度加倍将导致运行时增加大约8倍。
简单地说，在乘数函数mul（）中，循环的顺序可以任意改变。结果表明，它对运行时有很大的影响：最慢的运行时变量与最快的运行时变量之比约为1.73。
本文只描述最快的变量；本文末尾的附加代码（matr_mul dim）中显示了测试变量的其余部分。MQ5文件）。在这方面，OpenCL编程指南（AAFTAB Munshi、Benedict R.Gaster、Timothy G.Mattson、James Fung、Dan Ginsburg）中作了以下陈述（第512页）：
显然，这些并不是我们可以为最初的“非并行”代码实现的所有优化。有些与硬件（（S）SSEX指令相关，而另一些则是纯算法，如Strassen算法、铜匠Vinogratt算法等。注：Strassen算法的乘法矩阵大小非常小，只有6464，这导致了比传统算法更快的速度。在本文中，我们将研究线性尺寸达几千（约5000）的快速乘法矩阵。
2.2。
在OpenCL中的首次实现
现在，我们将该算法迁移到opencl以创建rows1（row 1）*cols2（column 2）线程，这将从内核中删除两个外部循环。每个线程执行colsrow迭代，以便内部循环保持内核的一部分。
由于我们必须为OpenCL内核创建三个线性缓冲区，因此有必要重新调整初始算法，使其尽可能接近内部会计方法。在带有线性缓冲区的“单核CPU”上，它与内核代码一起提供“非并行”程序代码。二维数组代码的最佳性并不意味着它的模拟对于线性缓冲区也是最佳的：所有的测试都必须重复。因此，我们再次选择C-R-C R作为线性代数中矩阵乘法标准逻辑对应的初始变量。
也就是说，为了避免矩阵/缓冲区寻址混乱，回答主要问题：如果矩阵matr（m行乘以n列）在GPU内存中被安排为线性缓冲区，我们如何计算元素matr[行][列]的线性偏移量？
事实上，GPU内存中的矩阵布局没有固定的顺序，因为它仅由问题的逻辑决定。例如，两个矩阵的元素在缓冲区排列上可能不同，因为矩阵在涉及的矩阵乘法算法方面是不对称的，即第一个矩阵的行乘以第二个矩阵的列。在每次内核迭代中，这种重新排列会极大地影响从全局GPU内存顺序读取矩阵元素的性能。
算法的第一个实现将具有相同的矩阵数组-按行一阶排列。首先，将第一行元素放入缓冲区，然后将第二行中的所有元素放入缓冲区，依此类推。下面是一个平面公式，线性存储上方矩阵matr[m（rows）][n（columns）的二维表示：
3
图13。GPU缓冲区中二维索引空间转换为线性置换矩阵的算法
上图还说明了线性存储中列优先级的二维矩阵表示。下面的
是我们首先在OpenCL设备上实现的一个代码，有少量删除：
//+------------------------------------------------------------------+ //| matr_mul_1dim.mq5 | //+------------------------------------------------------------------+ #property script_show_inputs #define ROWS1 2000 // rows in the first matrix #define COLSROWS 2000 // columns in the first matrix = rows in the second matrix #define COLS2 2000 // columns in the second matrix #define REALTYPE float REALTYPE first[]; // first linear buffer (matrix) rows1 * colsrows REALTYPE second[]; // second buffer colsrows * cols2 REALTYPE thirdGPU[ ]; // product - also a buffer rows1 * cols2 REALTYPE thirdCPU[ ]; // product - also a buffer rows1 * cols2 input int _device=1; // here is the device; it can be changed (now 4870) string d2s(double arg,int dig) { return DoubleToString(arg,dig); } string i2s(long arg) { return IntegerToString(arg); } //+------------------------------------------------------------------+ const string clSrc= "#define COLS2 "+i2s(COLS2)+" /r/n" "#define COLSROWS "+i2s(COLSROWS)+" /r/n" "#define REALTYPE float /r/n" " /r/n" "__kernel void matricesMul( __global REALTYPE *in1, /r/n" " __global REALTYPE *in2, /r/n" " __global REALTYPE *out ) /r/n" "{ /r/n" " int r = get_global_id( 0 ); /r/n" " int c = get_global_id( 1 ); /r/n" " for( int cr = 0; cr < COLSROWS; cr ++ ) /r/n" " out[ r * COLS2 + c ] += /r/n" " in1[ r * COLSROWS + cr ] * in2[ cr * COLS2 + c ]; /r/n" "} /r/n"; //+------------------------------------------------------------------+ //| Main matrix multiplication function; | //| Input matrices are already generated, | //| the output matrix is initialized to zeros | //+------------------------------------------------------------------+ void mulCPUOneCore() { //--- c-r-cr: 11.544 s //st = GetTickCount( ); for(int c=0; c<COLS2; c++) for(int r=0; r<ROWS1; r++) for(int cr=0; cr<COLSROWS; cr++) thirdCPU[r*COLS2+c]+=first[r*COLSROWS+cr]*second[cr*COLS2+c]; return; } //+------------------------------------------------------------------+ //| Script program start function | //+------------------------------------------------------------------+ void OnStart() { initAllDataCPU(); //--- start working with non-parallel version ("bare" CPU, single core) //--- calculate the output matrix on a single core CPU uint st=GetTickCount(); mulCPUOneCore(); //--- output total calculation time double timeCPU=(GetTickCount()-st)/1000.; Print("CPUTime = "+d2s(timeCPU,3)); //--- start working with OCL int clCtx; // context handle int clPrg; // handle to the program on the device int clKrn; // kernel handle int clMemIn1; // first (input) buffer handle int clMemIn2; // second (input) buffer handle int clMemOut; // third (output) buffer handle //--- start calculating the program runtime on GPU //st = GetTickCount( ); initAllDataGPU(clCtx,clPrg,clKrn,clMemIn1,clMemIn2,clMemOut); //--- start calculating total OCL code runtime st=GetTickCount(); executeGPU(clKrn); //--- create a buffer for reading and read the result; we will need it later REALTYPE buf[]; readOutBuf(clMemOut,buf); //--- stop calculating the total program runtime //--- together with the time required for retrieval of data from GPU and transferring it back to RAM double timeGPUTotal=(GetTickCount()-st)/1000.; Print("OpenCL total: time = "+d2s(timeGPUTotal,3)+" sec."); destroyOpenCL(clCtx,clPrg,clKrn,clMemIn1,clMemIn2,clMemOut); //--- calculate the time elapsed Print("CPUTime / GPUTotalTime = "+d2s(timeCPU/timeGPUTotal,3)); //--- debugging: random checks. Multiplication accuracy is checked directly //--- on the initial and output matrices using a few dozen examples for(int i=0; i<10; i++) checkRandom(buf,ROWS1,COLS2); Print("________________________"); return; } //+------------------------------------------------------------------+ //| initAllDataCPU | //+------------------------------------------------------------------+ void initAllDataCPU() { //--- initialize random number generator MathSrand(( int) TimeLocal()); Print("======================================="); Print("1st OCL martices mul: device = "+i2s(_device)+"; ROWS1 = " +i2s(ROWS1)+ "; COLSROWS = "+i2s(COLSROWS)+"; COLS2 = "+i2s(COLS2)); //--- set the required sizes of linear representations of the input and output matrices ArrayResize(first,ROWS1*COLSROWS); ArrayResize(second,COLSROWS*COLS2); ArrayResize(thirdGPU,ROWS1*COLS2); ArrayResize(thirdCPU,ROWS1*COLS2); //--- generate both input matrices and initialize the output to zeros genMatrices(); ArrayInitialize( thirdCPU, 0.0 ); ArrayInitialize( thirdGPU, 0.0 ); return; } //+------------------------------------------------------------------+ //| initAllDataCPU | //| lay out in row-major order, Matr[ M (rows) ][ N (columns) ]: | //| Matr[row][column] = buff[row * N(columns in the matrix) + column]| //| generate initial matrices; this generation is not reflected | //| in the final runtime calculation | //| buffers are filled in row-major order! | //+------------------------------------------------------------------+ void genMatrices() { for(int r=0; r<ROWS1; r++) for(int c=0; c<COLSROWS; c++) first[r*COLSROWS+c]=genVal(); for(int r=0; r<COLSROWS; r++) for(int c=0; c<COLS2; c++) second[r*COLS2+c]=genVal(); return; } //+------------------------------------------------------------------+ //| genVal | //| generate one value of the matrix element: | //| uniformly distributed value lying in the range [-0.5; 0.5] | //+------------------------------------------------------------------+ REALTYPE genVal() { return(REALTYPE)((MathRand()-16383.5)/32767.); } //+------------------------------------------------------------------+ //| initAllDataGPU | //+------------------------------------------------------------------+ void initAllDataGPU(int &clCtx, // context int& clPrg, // program on the device int& clKrn, // kernel int& clMemIn1, // first (input) buffer int& clMemIn2, // second (input) buffer int& clMemOut) // third (output) buffer { //--- write the kernel code to a file WriteCLProgram(); //--- create context, program and kernel clCtx = CLContextCreate( _device ); clPrg = CLProgramCreate( clCtx, clSrc ); clKrn = CLKernelCreate( clPrg, "matricesMul" ); //--- create all three buffers for the three matrices //--- first matrix - input clMemIn1=CLBufferCreate(clCtx,ROWS1 *COLSROWS*sizeof(REALTYPE),CL_MEM_READ_WRITE); //--- second matrix - input clMemIn2=CLBufferCreate(clCtx,COLSROWS*COLS2 *sizeof(REALTYPE),CL_MEM_READ_WRITE); //--- third matrix - output clMemOut=CLBufferCreate(clCtx,ROWS1 *COLS2 *sizeof(REALTYPE),CL_MEM_READ_WRITE); //--- set arguments to the kernel CLSetKernelArgMem(clKrn,0,clMemIn1); CLSetKernelArgMem(clKrn,1,clMemIn2); CLSetKernelArgMem(clKrn,2,clMemOut); //--- write the generated matrices to the device buffers CLBufferWrite(clMemIn1,first); CLBufferWrite(clMemIn2,second); CLBufferWrite(clMemOut,thirdGPU); // 0.0 everywhere return; } //+------------------------------------------------------------------+ //| WriteCLProgram | //+------------------------------------------------------------------+ void WriteCLProgram() { int h=FileOpen("matr_mul_OCL_1st.cl",FILE_WRITE|FILE_TXT|FILE_ANSI); FileWrite(h,clSrc); FileClose(h); } //+------------------------------------------------------------------+ //| executeGPU | //+------------------------------------------------------------------+ void executeGPU(int clKrn) { //--- set the workspace parameters for the task and execute the OpenCL program uint offs[ 2 ] = { 0, 0 }; uint works[ 2 ] = { ROWS1, COLS2 }; bool ex=CLExecute(clKrn,2,offs,works); return; } //+------------------------------------------------------------------+ //| readOutBuf | //+------------------------------------------------------------------+ void readOutBuf(int clMemOut,REALTYPE &buf[]) { ArrayResize(buf,COLS2*ROWS1); //--- buf - a copy of what is written to the buffer thirdGPU[] uint read=CLBufferRead(clMemOut,buf); Print("read = "+i2s(read)+" elements"); return; } //+------------------------------------------------------------------+ //| destroyOpenCL | //+------------------------------------------------------------------+ void destroyOpenCL(int clCtx,int clPrg,int clKrn,int clMemIn1,int clMemIn2,int clMemOut) { //--- destroy all that was created for calculations on the OpenCL device in reverse order CLBufferFree(clMemIn1); CLBufferFree(clMemIn2); CLBufferFree(clMemOut); CLKernelFree(clKrn); CLProgramFree(clPrg); CLContextFree(clCtx); return; } //+------------------------------------------------------------------+ //| checkRandom | //| random check of calculation accuracy | //+------------------------------------------------------------------+ void checkRandom(REALTYPE &buf[],int rows,int cols) { int r0 = genRnd( rows ); int c0 = genRnd( cols ); REALTYPE sum=0.0; for(int runningIdx=0; runningIdx<COLSROWS; runningIdx++) sum+=first[r0*COLSROWS+runningIdx]* second[runningIdx*COLS2+c0]; //--- element of the buffer m[] REALTYPE bufElement=buf[r0*COLS2+c0]; //--- element of the matrix not calculated in OpenCL REALTYPE CPUElement=thirdCPU[r0*COLS2+c0]; Print("sum( "+i2s(r0)+","+i2s(c0)+" ) = "+d2s(sum,8)+ "; thirdCPU[ "+i2s(r0)+","+i2s(c0)+" ] = "+d2s(CPUElement,8)+ "; buf[ "+i2s(r0)+","+i2s(c0)+" ] = "+d2s(bufElement,8)); return; } //+------------------------------------------------------------------+ //| genRnd | //+------------------------------------------------------------------+ int genRnd(int max) { return(int)(MathRand()/32767.*max); }
清单2。opencl程序中
的第一个实现
最后两个函数用于验证计算的准确性。完整的代码（matr_mul_1dim.mq5）附在本文末尾。注：尺寸不必只对应于正方形矩阵。
进一步的更改几乎总是关于内核代码的，所以我将在下面讨论内核修改代码。
为了方便计算精度由浮点型转换为双精度型，引入了realtype类型。应该注意的是，realtype类型不仅在宿主程序中声明，而且在内核中声明。如有必要，这种类型的任何更改都应该同时在两个地方进行——主机程序的定义和内核代码的定义。
代码执行结果（以下所有浮点数据类型）
CPU（opencl，_device=0）：
2012.05.20 22:14:57 matr_mul_1dim (EURUSD,H1) CPUTime / GPUTotalTime = 12.479 2012.05.20 22:14:57 matr_mul_1dim (EURUSD,H1) OpenCL total: time = 9.266 sec. 2012.05.20 22:14:57 matr_mul_1dim (EURUSD,H1) read = 4000000 elements 2012.05.20 22:14:48 matr_mul_1dim (EURUSD,H1) CPUTime = 115.628 2012.05.20 22:12:52 matr_mul_1dim (EURUSD,H1) 1st OCL martices mul: device = 0; ROWS1 = 2000; COLSROWS = 2000; COLS2 = 2000 2012.05.20 22:12:52 matr_mul_1dim (EURUSD,H1) =======================================
例如，在Radeon HD 4870（_device=1）上：
2012.05.27 01:40:50 matr_mul_1dim (EURUSD,H1) CPUTime / GPUTotalTime = 9.002 2012.05.27 01:40:50 matr_mul_1dim (EURUSD,H1) OpenCL total: time = 12.729 sec. 2012.05.27 01:40:50 matr_mul_1dim (EURUSD,H1) read = 4000000 elements 2012.05.27 01:40:37 matr_mul_1dim (EURUSD,H1) CPUTime = 114.583 2012.05.27 01:38:42 matr_mul_1dim (EURUSD,H1) 1st OCL martices mul: device = 1; ROWS1 = 2000; COLSROWS = 2000; COLS2 = 2000 2012.05.27 01:38:42 matr_mul_1dim (EURUSD,H1) =======================================
如您所见，这个内核在GPU上的实现要慢得多。然而，我们还没有特别优化GPU。
几个结论：
- 将矩阵形式从二维改为线性（在相应设备上执行的程序形式）对程序连续版本的总体运行时间没有显著影响。
- 选择线性代数中最直观的矩阵乘法定义匹配算法作为初始变量进行进一步优化。在某种程度上，它比最快的算法慢，但从未来GPU加速的角度来看，这个因素并不重要。
- 只有在将缓冲区读取到RAM（而不是clexecute（）命令）后才计算运行时。根据metadriver，其背后的原因可能是：metadriver:clbufferread（）在从缓冲区读取程序之前等待程序实际完成。clexecute（）实际上是一个异步队列函数。在cl代码用完之前，它会立即返回结果。
- GPU计算向导通常不计算内核的运行时，而是计算与各种对象（存储、算术等）相关联的吞吐量功能。我们可以并将重复下面的步骤。
如我们所知，一个维度为2000的矩阵需要每个元素2*2000的加法/乘法。将矩阵元素数（2000*2000）相乘，可以得到160亿个浮点数据运算。也就是说，在CPU上执行此操作需要115.628秒，对应于以下数据流速度
另一方面，
也知道，到目前为止，矩阵维2000的“单核CPU”的最快运行时间只有83.663秒（见我们第一个没有OpenCL的代码）。
我们使用这些数据作为参考，并作为优化的起点。
与之类似，CPU 上利用 OpenCL 进行的计算会得到：
最后，
计算GPU的吞吐量：
2.3。消除不连贯的数据访问
从这个内核代码中，您可以很容易地看到几个非最佳项。
看看内核中的循环体：
for( int cr = 0; cr < COLSROWS; cr ++ ) out[ r * COLS2 + c ] += in1[ r * COLSROWS + cr ] * in2[ cr * COLS2 + c ];
很容易看出，当运行循环计数器（CR++）时，连续数据是从1[]中的第一个缓冲区获取的。以cols2的“间隔”检索来自第二个缓冲区in2[]的数据。换句话说，从第二个缓冲区中检索到的数据的主要部分是无用的，因为存储请求将不连贯（见1.2.1）。合并内存请求）。针对这种情况，通过改变数组in2[]中索引的计算公式及其生成方式，可以修改三个代码：
– 内核代码：
for( int cr = 0; cr < COLSROWS; cr ++ ) out[ r * COLS2 + c ] += in1[ r * COLSROWS + cr ] * in2[ cr + c * COLSROWS ];
现在，当循环计数器 (cr++) 值变化时，两个数组的数据都会被顺序获取，不带任何“间隔”。
– genMatrices() 中的缓冲区填充代码。现在，应按列优先顺序（而不是开头处使用的行优先顺序）将其填入：
for( int r = 0; r < COLSROWS; r ++ ) for( int c = 0; c < COLS2; c ++ ) /// second[ r * COLS2 + c ] = genVal( ); second[ r + c * COLSROWS ] = genVal( );
– checkRandom() 函数中的验证码：
for( int runningIdx = 0; runningIdx < COLSROWS; runningIdx ++ ) ///sum += first[ r0 * COLSROWS + runningIdx ] * second[ runningIdx * COLS2 + c0 ]; sum += first[ r0 * COLSROWS + runningIdx ] * second[ runningIdx + c0 * COLSROWS ];
CPU 上的性能结果：
2012.05.24 02:59:22 matr_mul_1dim_coalesced (EURUSD,H1) CPUTime / GPUTotalTime = 16.207 2012.05.24 02:59:22 matr_mul_1dim_coalesced (EURUSD,H1) OpenCL total: time = 5.756 sec. 2012.05.24 02:59:22 matr_mul_1dim_coalesced (EURUSD,H1) read = 4000000 elements 2012.05.24 02:59:16 matr_mul_1dim_coalesced (EURUSD,H1) CPUTime = 93.289 2012.05.24 02:57:43 matr_mul_1dim_coalesced (EURUSD,H1) 1st OCL martices mul: device = 0; ROWS1 = 2000; COLSROWS = 2000; COLS2 = 2000 2012.05.24 02:57:43 matr_mul_1dim_coalesced (EURUSD,H1) =======================================
Radeon HD 4870：
2012.05.27 01:50:43 matr_mul_1dim_coalesced (EURUSD,H1) CPUTime / GPUTotalTime = 7.176 2012.05.27 01:50:43 matr_mul_1dim_coalesced (EURUSD,H1) OpenCL total: time = 12.979 sec. 2012.05.27 01:50:43 matr_mul_1dim_coalesced (EURUSD,H1) read = 4000000 elements 2012.05.27 01:50:30 matr_mul_1dim_coalesced (EURUSD,H1) CPUTime = 93.133 2012.05.27 01:48:57 matr_mul_1dim_coalesced (EURUSD,H1) 1st OCL martices mul: device = 1; ROWS1 = 2000; COLSROWS = 2000; COLS2 = 2000 2012.05.27 01:48:57 matr_mul_1dim_coalesced (EURUSD,H1) =======================================
可以看出，对数据的一致访问对GPU上的运行时几乎没有影响，但它显著提高了CPU上的运行时。这很可能与后期优化的因素有关，特别是访问全局变量的高延迟，我们需要尽快消除这些因素。
throughput_arithmetic_CPU_OCL = 16 000000000 / 5.756 ~ 2.780 GFlops. throughput_arithmetic_GPU_OCL = 16 000000000 / 12.979 ~ 1.233 GFlops.
有关这个新的内核代码，请参阅本文末尾的matr_mul_1dim_coalesced.mq5。
内核代码如下：
const string clSrc = "#define COLS2 " + i2s( COLS2 ) + " /r/n" "#define COLSROWS " + i2s( COLSROWS ) + " /r/n" "#define REALTYPE float /r/n" " /r/n" "__kernel void matricesMul( __global REALTYPE *in1, /r/n" " __global REALTYPE *in2, /r/n" " __global REALTYPE *out ) /r/n" "{ /r/n" " int r = get_global_id( 0 ); /r/n" " int c = get_global_id( 1 ); /r/n" " for( int cr = 0; cr < COLSROWS; cr ++ ) /r/n" " out[ r * COLS2 + c ] += /r/n" " in1[ r * COLSROWS + cr ] * in2[ cr + c * COLSROWS ]; /r/n" "} /r/n";
清单3。内核
，具有合并的全局存储数据访问
现在，我们继续进一步优化。
2.4. 移除输出矩阵中“代价高昂”的全局 GPU 存储访问
众所周知，全球GPU存储访问的延迟非常高（约600-800个周期）。例如，一次执行两个数字的加法会导致大约20个周期的延迟。在GPU上进行计算时，优化的主要目标是通过增加计算吞吐量来隐藏延迟。在先前开发的内核循环中，我们不断地访问全局存储元素，从而浪费了大量时间。
现在，我们引入局部变量和内核（由于它是工作寄存器单元内核中的一个私有变量，可以更快地访问多次），当循环结束时，我们将获得的和和和值逐个分配给输出数据的元素：
const string clSrc = "#define COLS2 " + i2s( COLS2 ) + " /r/n" "#define COLSROWS " + i2s( COLSROWS ) + " /r/n" "#define REALTYPE float /r/n" " /r/n" "__kernel void matricesMul( __global REALTYPE *in1, /r/n" " __global REALTYPE *in2, /r/n" " __global REALTYPE *out ) /r/n" "{ /r/n" " int r = get_global_id( 0 ); /r/n" " int c = get_global_id( 1 ); /r/n" " REALTYPE sum = 0.0; /r/n" " for( int cr = 0; cr < COLSROWS; cr ++ ) /r/n" " sum += in1[ r * COLSROWS + cr ] * in2[ cr + c * COLSROWS ]; /r/n" " out[ r * COLS2 + c ] = sum; /r/n" "} /r/n" ;
清单4。引入私有变量计算无向积计算循环中的累积和
有关完整的源代码文件，请参阅本文末尾附带的matr_mul_sum_local.mq5。
CPU:
2012.05.24 03:28:17 matr_mul_sum_local (EURUSD,H1) CPUTime / GPUTotalTime = 24.863 2012.05.24 03:28:16 matr_mul_sum_local (EURUSD,H1) OpenCL total: time = 3.759 sec. 2012.05.24 03:28:16 matr_mul_sum_local (EURUSD,H1) read = 4000000 elements 2012.05.24 03:28:12 matr_mul_sum_local (EURUSD,H1) CPUTime = 93.460 2012.05.24 03:26:39 matr_mul_sum_local (EURUSD,H1) 1st OCL martices mul: device = 0; ROWS1 = 2000; COLSROWS = 2000; COLS2 = 2000
GPU HD 4870：
2012.05.27 01:57:30 matr_mul_sum_local (EURUSD,H1) CPUTime / GPUTotalTime = 69.541 2012.05.27 01:57:30 matr_mul_sum_local (EURUSD,H1) OpenCL total: time = 1.326 sec. 2012.05.27 01:57:30 matr_mul_sum_local (EURUSD,H1) read = 4000000 elements 2012.05.27 01:57:28 matr_mul_sum_local (EURUSD,H1) CPUTime = 92.212 2012.05.27 01:55:56 matr_mul_sum_local (EURUSD,H1) 1st OCL martices mul: device = 1; ROWS1 = 2000; COLSROWS = 2000; COLS2 = 2000 2012.05.27 01:55:56 matr_mul_sum_local (EURUSD,H1) =======================================
这是一次真正的生产率提升！
throughput_arithmetic_CPU_OCL = 16 000000000 / 3.759 ~ 4.257 GFlops. throughput_arithmetic_GPU_OCL = 16 000000000 / 1.326 ~ 12.066 GFlops.
在顺序优化过程中，我们努力遵循的主要原则是：首先，您需要以最完美的方式重新排列数据结构，使其适合于给定的任务，特别是基本的硬件。然后，采用MAD（）或FMA（）等快速算法进行优化。请记住，顺序优化不一定会导致性能改进——这是不保证的。
2.5。改进内核执行的操作
在并行编程中，根据并行操作结构，将组织计算的间接成本（花费时间）降到最低是非常重要的。在Dimension 2000的矩阵中，计算输出矩阵元素的工作单元将执行总任务的1/4000000工作量。
这显然太多了，与在硬件上执行计算的实际单元数相差甚远。现在，在这个新版本的内核中，我们将计算整个矩阵行，而不是一个元素。
重要的是，现在任务空间已经从二维变为一维，成为一个全维的整行（而不是矩阵的单个元素），并且已经在内核的每个任务中进行了计算。因此，任务空间被转换为矩阵的行数。
4
图14。输出矩阵整行计算示意图
内核代码变得更加复杂：
const string clSrc = "#define COLS2 " + i2s( COLS2 ) + " /r/n" "#define COLSROWS " + i2s( COLSROWS ) + " /r/n" "#define REALTYPE float /r/n" " /r/n" "__kernel void matricesMul( __global REALTYPE *in1, /r/n" " __global REALTYPE *in2, /r/n" " __global REALTYPE *out ) /r/n" "{ /r/n" " int r = get_global_id( 0 ); /r/n" " REALTYPE sum; /r/n" " for( int c = 0; c < COLS2; c ++ ) /r/n" " { /r/n" " sum = 0.0; /r/n" " for( int cr = 0; cr < COLSROWS; cr ++ ) /r/n" " sum += in1[ r * COLSROWS + cr ] * in2[ cr + c * COLSROWS ]; /r/n" " out[ r * COLS2 + c ] = sum; /r/n" " } /r/n" "} /r/n" ;
清单5。用于输出矩阵整行计算的内核
此外，任务维数亦已于 executeGPU( ) 函数中更改：
void executeGPU( int clKrn ) { //--- set parameters of the task workspace and execute the OpenCL program uint offs[ 1 ] = { 0 }; uint works[ 1 ] = { ROWS1 }; bool ex = CLExecute( clKrn, 1, offs, works ); return; }
性能结果（完整源代码见Matr_Mul_Row_Calc.mq5）：
CPU:
2012.05.24 15:56:24 matr_mul_row_calc (EURUSD,H1) CPUTime / GPUTotalTime = 17.385 2012.05.24 15:56:24 matr_mul_row_calc (EURUSD,H1) OpenCL total: time = 5.366 sec. 2012.05.24 15:56:24 matr_mul_row_calc (EURUSD,H1) read = 4000000 elements 2012.05.24 15:56:19 matr_mul_row_calc (EURUSD,H1) CPUTime = 93.288 2012.05.24 15:54:45 matr_mul_row_calc (EURUSD,H1) 1st OCL martices mul: device = 0; ROWS1 = 2000; COLSROWS = 2000; COLS2 = 2000 2012.05.24 15:54:45 matr_mul_row_calc (EURUSD,H1) =======================================
GPU 4870：
2012.05.27 02:24:10 matr_mul_row_calc (EURUSD,H1) CPUTime / GPUTotalTime = 55.119 2012.05.27 02:24:10 matr_mul_row_calc (EURUSD,H1) OpenCL total: time = 1.669 sec. 2012.05.27 02:24:10 matr_mul_row_calc (EURUSD,H1) read = 4000000 elements 2012.05.27 02:24:08 matr_mul_row_calc (EURUSD,H1) CPUTime = 91.994 2012.05.27 02:22:35 matr_mul_row_calc (EURUSD,H1) 1st OCL martices mul: device = 1; ROWS1 = 2000; COLSROWS = 2000; COLS2 = 2000 2012.05.27 02:22:35 matr_mul_row_calc (EURUSD,H1) =======================================
我们可以看到，CPU上的运行时间明显恶化，而GPU上的运行时间稍差（尽管不明显）。并非所有这些都是那么糟糕：这暂时恶化了本地的战略格局，只是为了进一步大幅提高其绩效。
throughput_arithmetic_CPU_OCL = 16 000000000 / 5.366 ~ 2.982 GFlops. throughput_arithmetic_GPU_OCL = 16 000000000 / 1.669 ~ 9.587 GFlops.
2.6。将第一个数组的行迁移到专用内存
矩阵乘法算法的主要特点是具有大量伴随结果的乘法。对该算法进行适当而高质量的优化意味着数据传输的最小化。但到目前为止，在计算无向累积积主环的过程中，我们所有的核修改都存储在全局内存的三个矩阵中的两个。
这意味着每个无向产品（实际上是每个输出矩阵元素）的所有输入数据都会在整个存储层次结构（从全局到私有）中通过相关延迟进行连续优化。对于输出矩阵的每一个计算行，通过确保每个工作单元重用第一个矩阵的同一行来减少流量。
5
图15。将第一个矩阵的行迁移到工作单元的专用内存
不需要对主机程序代码进行任何更改。此外，内核中的变化非常微妙。考虑到中间的一维私有数组是在内核中生成的，GPU试图将它放在执行内核的单元的私有内存中。将第一个矩阵所需的行从全局内存复制到私有内存。换句话说，值得注意的是，即使是这种复制也相对较快。关键是将第一个数组行元素从全局内存复制到私有内存的成本是一致性方法，与计算输出矩阵行的主双精度循环相比，复制的间接成本相当适中。
内核代码（主循环中注释掉的代码是上一版本中已经存在的代码）：
const string clSrc = "#define COLS2 " + i2s( COLS2 ) + " /r/n" "#define COLSROWS " + i2s( COLSROWS ) + " /r/n" "#define REALTYPE float /r/n" " /r/n" "__kernel void matricesMul( __global REALTYPE *in1, /r/n" " __global REALTYPE *in2, /r/n" " __global REALTYPE *out ) /r/n" "{ /r/n" " int r = get_global_id( 0 ); /r/n" " REALTYPE rowbuf[ COLSROWS ]; /r/n" " for( int col = 0; col < COLSROWS; col ++ ) /r/n" " rowbuf[ col ] = in1[ r * COLSROWS + col ]; /r/n" " REALTYPE sum; /r/n" " /r/n" " for( int c = 0; c < COLS2; c ++ ) /r/n" " { /r/n" " sum = 0.0; /r/n" " for( int cr = 0; cr < COLSROWS; cr ++ ) /r/n" " ///sum += in1[ r * COLSROWS + cr ] * in2[ cr + c * COLSROWS ]; /r/n" " sum += rowbuf[ cr ] * in2[ cr + c * COLSROWS ]; /r/n" " out[ r * COLS2 + c ] = sum; /r/n" " } /r/n" "} /r/n" ;
清单6。内核与工作单元的私有内存中的第一个矩阵行。
CPU：
2012.05.27 00:51:46 matr_mul_row_in_private (EURUSD,H1) CPUTime / GPUTotalTime = 18.587 2012.05.27 00:51:46 matr_mul_row_in_private (EURUSD,H1) OpenCL total: time = 4.961 sec. 2012.05.27 00:51:46 matr_mul_row_in_private (EURUSD,H1) read = 4000000 elements 2012.05.27 00:51:41 matr_mul_row_in_private (EURUSD,H1) CPUTime = 92.212 2012.05.27 00:50:08 matr_mul_row_in_private (EURUSD,H1) 1st OCL martices mul: device = 0; ROWS1 = 2000; COLSROWS = 2000; COLS2 = 2000 2012.05.27 00:50:08 matr_mul_row_in_private (EURUSD,H1) =======================================
GPU：
2012.05.27 02:28:49 matr_mul_row_in_private (EURUSD,H1) CPUTime / GPUTotalTime = 69.242 2012.05.27 02:28:49 matr_mul_row_in_private (EURUSD,H1) OpenCL total: time = 1.327 sec. 2012.05.27 02:28:49 matr_mul_row_in_private (EURUSD,H1) read = 4000000 elements 2012.05.27 02:28:47 matr_mul_row_in_private (EURUSD,H1) CPUTime = 91.884 2012.05.27 02:27:15 matr_mul_row_in_private (EURUSD,H1) 1st OCL martices mul: device = 1; ROWS1 = 2000; COLSROWS = 2000; COLS2 = 2000 2012.05.27 02:27:15 matr_mul_row_in_private (EURUSD,H1) =======================================
throughput_arithmetic_CPU_OCL = 16 000000000 / 4.961 ~ 3.225 GFlops. throughput_arithmetic_GPU_OCL = 16 000000000 / 1.327 ~ 12.057 GFlops.
CPU吞吐量与前一个保持大致相同的水平，而GPU吞吐量恢复到有史以来的最高水平，但仍处于新的容量中。注意：CPU吞吐量似乎当场冻结，但略有不稳定；GPU吞吐量飙升（尽管并非总是如此）。
请注意，实际的算术吞吐量稍高，因为第一个矩阵的行被复制到私有内存中，所以现在执行的操作比以前多。但是，它对最终吞吐量的估计几乎没有影响。
有关源代码，请参见u private.mq5中的matr_mul_row_。
2.7。将第二个数组的列迁移到本地内存
现在，很容易猜测下一步该做什么。我们已经采取措施隐藏与输出和第一个输入矩阵相关的延迟。第二个矩阵在左边。
对矩阵乘法中使用的无向积进行了更详细的研究，结果表明，在计算输出矩阵行的过程中，该组中的所有单元都通过该设备从第二乘法矩阵的同一列重新优化数据。如下图所示：
6
图16。将第二个矩阵的列迁移到工作组的本地数据共享
如果构成工作组的工作单元在输出矩阵的行开始计算之前将第二个矩阵的列复制到工作组内存，则从全局内存迁移的间接成本也将降低。
这需要更改内核和主机程序。最重要的改变是为每个内核设置本地内存。它应该是显式的，因为OpenCL不支持动态内存分配。因此，应该首先将足够大的存储对象放在主机中，以便在内核中进行进一步的处理。
只有这样，工作单元才能在执行内核时将第二个矩阵的列复制到本地内存。这是通过利用工作组所有工作单元的迭代循环分布来并行完成的。但是，所有复制必须在工作单元开始其主操作（计算输出矩阵行）之前完成。
因此，还应在负责复制的循环之后插入以下命令：
barrier(CLK_LOCAL_MEM_FENCE);
这是一个“本地内存屏障”，确保组中的工作单元可以“看到”与其他单元协调的特定状态下的本地内存。工作组中的所有工作单元必须首先执行所有到屏障的命令，然后才能继续执行内核。换言之，这一障碍是工作组各工作单位之间的一种特殊同步机制。
OpenCL中的工作组之间没有同步机制。
下图说明了活动的障碍：
7
图17。主动屏障图
实际上，似乎只有工作组中的工作单元严格地同时执行代码。这只是OpenCL编程模型的一个抽象。
到目前为止，我们在不同的工作单元上执行的内核代码不需要同步操作，因为它们之间没有将在内核中编程的显式通信；而且，甚至不需要同步操作。但是，在这个内核中需要同步，因为本地数组的填充是并行分布在工作组的所有单元中的。
换句话说，每个工作单元将其值写入本地数据共享（这里是一个数组），而不知道其他工作单元在写入过程中的距离。设置一个屏障，使特定的工作单元在必要时（即，在完全生成本地数组之前）才继续执行内核。
您需要了解，这种优化几乎不会对CPU性能产生任何有益的影响：根据Intel的OpenCL优化指南，当在CPU上执行某个内核时，所有OpenCL内存对象都由硬件缓存，因此使用本地内存的显式缓存只会增加不必要（但中等）的间接成本。
这是另一个值得一提的观点——作者花了很多时间。它关注的是局部变量不能在内核函数的头中传递（即，在编译阶段，在终端开发人员的构建的当前实现中）。这背后的原因是，为了将内存分配为一个与内核函数无关的变量，我们必须首先使用clbuffercreate（）函数在CPU内存中显式创建该对象，并将其大小显式指定为一个函数参数。此函数返回一个内存对象句柄，该句柄进一步存储在全局GPU内存中，因为这是它唯一的庇护所。
但是，本地内存与全局内存的类型不同，因此创建的内存对象不能放在工作组的本地内存中。
功能齐全的opencl api允许显式分配所需大小的内存，指针NULL指向内核的独立变量，即使不是（clsetkernelarg（）函数创建内存对象）。但是，将clsetkernelargmem（）函数模拟为功能齐全的API函数的mql5不允许我们在不创建它的情况下将分配给独立变量的内存大小传递给内存对象本身。我们可以传递给clsetkernelargmem（）函数的只是在全局CPU中生成并设计为迁移到全局GPU内存的缓冲区句柄。这是一个悖论。
幸运的是，在这个内核中有一种使用本地缓冲区的等效方法。只需使用_local修饰符在内核主体中声明缓冲区。在这种情况下，分配给这个工作组的本地内存是在运行时确定的，而不是在编译期间。
遵循内核中继屏障的命令（代码中的屏障线用红色标记）与以前的优化过程基本相同。主机程序代码仍然相同（有关源代码，请参阅matr_mul_col_local.mq5）。
新的内核代码如下：
const string clSrc = "#define COLS2 " + i2s( COLS2 ) + " /r/n" "#define COLSROWS " + i2s( COLSROWS ) + " /r/n" "#define REALTYPE float /r/n" " /r/n" "__kernel void matricesMul( __global REALTYPE *in1, /r/n" " __global REALTYPE *in2, /r/n" " __global REALTYPE *out ) /r/n" "{ /r/n" " int r = get_global_id( 0 ); /r/n" " REALTYPE rowbuf[ COLSROWS ]; /r/n" " for( int col = 0; col < COLSROWS; col ++ ) /r/n" " rowbuf[ col ] = in1[ r * COLSROWS + col ]; /r/n" " /r/n" " int idlocal = get_local_id( 0 ); /r/n" " int nlocal = get_local_size( 0 ); /r/n" " __local REALTYPE colbuf[ COLSROWS ] ; /r/n" " /r/n" " REALTYPE sum; /r/n" " for( int c = 0; c < COLS2; c ++ ) /r/n" " { /r/n" " for( int cr = idlocal; cr < COLSROWS; cr = cr + nlocal ) /r/n" " colbuf[ cr ] = in2[ cr + c * COLSROWS ]; /r/n" " barrier( CLK_LOCAL_MEM_FENCE ); /r/n" " /r/n" " sum = 0.0; /r/n" " for( int cr = 0; cr < COLSROWS; cr ++ ) /r/n" " sum += rowbuf[ cr ] * colbuf[ cr ]; /r/n" " out[ r * COLS2 + c ] = sum; /r/n" " } /r/n" "} /r/n" ;
清单7。第二个数组的列
迁移到工作组本地内存
CPU：
2012.05.27 06:31:46 matr_mul_col_local (EURUSD,H1) CPUTime / GPUTotalTime = 17.630 2012.05.27 06:31:46 matr_mul_col_local (EURUSD,H1) OpenCL total: time = 5.227 sec. 2012.05.27 06:31:46 matr_mul_col_local (EURUSD,H1) read = 4000000 elements 2012.05.27 06:31:40 matr_mul_col_local (EURUSD,H1) CPUTime = 92.150 2012.05.27 06:30:08 matr_mul_col_local (EURUSD,H1) 1st OCL martices mul: device = 0; ROWS1 = 2000; COLSROWS = 2000; COLS2 = 2000 2012.05.27 06:30:08 matr_mul_col_local (EURUSD,H1) =======================================
GPU：
2012.05.27 06:21:36 matr_mul_col_local (EURUSD,H1) CPUTime / GPUTotalTime = 58.069 2012.05.27 06:21:36 matr_mul_col_local (EURUSD,H1) OpenCL total: time = 1.592 sec. 2012.05.27 06:21:36 matr_mul_col_local (EURUSD,H1) read = 4000000 elements 2012.05.27 06:21:34 matr_mul_col_local (EURUSD,H1) CPUTime = 92.446 2012.05.27 06:20:01 matr_mul_col_local (EURUSD,H1) 1st OCL martices mul: device = 1; ROWS1 = 2000; COLSROWS = 2000; COLS2 = 2000 2012.05.27 06:20:01 matr_mul_col_local (EURUSD,H1) =======================================
这两种情况都表现出性能下降，但并不严重。通过改变工作组的大小，可以提高（而不是降低）性能。最好将上面的示例用于不同的目的——演示如何使用本地内存对象。
有一个假设可以解释使用本地内存时性能下降的原因。大约两年前，habrahabr.ru发表了一篇题为“CUDA、GLSL和OpenMP与OpenCL（俄语）的比较”的论文。它说：
就是同一篇文章的下方，作者做出了如下论述：
换句话说，这是否意味着两年前发布的产品的本地内存不比全局内存快？上述出版时间表明，两年前，Radeon HD58xx系列图形卡已经上市，但据作者介绍，它们还远远不够完美。我很难相信这一点，尤其是当AMD推出了轰动的常青系列产品时。我想用更新后的显卡来检查，如HD69XX系列。
补充：打开GPU Caps Viewer（GPU性能查看器），您将在OpenCL选项卡中看到以下内容：
8
图18。hd 4870支持的opencl主参数
cl_设备本地内存类型：全局
语言规范（第41页表4.3）中对此参数的解释如下：
支持本地内存类型。它可以设置为cl_local，表示它是一个专用的本地内存存储，如sram或cl_global。
因此，HD4870本地内存实际上是全局内存的一部分，因此图形卡中的任何本地内存操作都不再有用，并且不会导致比全局内存条件出现更快的情况。这是另一个链接，包括AMD专家对高清4xx序列的澄清。这并不一定意味着你的图形卡有多坏；它只是告诉你在哪里可以找到与硬件相关的信息-在这种情况下，在GPU Caps Viewer中。
throughput_arithmetic_CPU_OCL = 16 000000000 / 5.227 ~ 3.061 GFlops. throughput_arithmetic_GPU_OCL = 16 000000000 / 1.592 ~ 10.050 GFlops.
最后，我们通过显式向量化内核来添加一些亮点。从第一个数组到私有内存阶段的行迁移派生的内核（matr_mul_row_in_private.mq5）将充当原始内核，因为它似乎是最快的。
2.8。核向量化
为了避免混淆，最好将操作分成几个阶段。在初始修改中，我们不会改变内核外部参数的数据类型，只会对内部循环的计算进行矢量化：
const string clSrc = "#define COLS2 " + i2s( COLS2 ) + " /r/n" "#define COLSROWS " + i2s( COLSROWS ) + " /r/n" "#define REALTYPE float /r/n" "#define REALTYPE4 float4 /r/n" " /r/n" "__kernel void matricesMul( __global REALTYPE *in1, /r/n" " __global REALTYPE *in2, /r/n" " __global REALTYPE *out ) /r/n" "{ /r/n" " int r = get_global_id( 0 ); /r/n" " REALTYPE rowbuf[ COLSROWS ]; /r/n" " for( int col = 0; col < COLSROWS; col ++ ) /r/n" " { /r/n" " rowbuf[ col ] = in1[r * COLSROWS + col ]; /r/n" " } /r/n" " /r/n" " REALTYPE sum; /r/n" " /r/n" " for( int c = 0; c < COLS2; c ++ ) /r/n" " { /r/n" " sum = 0.0; /r/n" " for( int cr = 0; cr < COLSROWS; cr += 4 ) /r/n" " sum += dot( ( REALTYPE4 ) ( rowbuf[ cr ], /r/n" " rowbuf[ cr + 1 ], /r/n" " rowbuf[ cr + 2 ], /r/n" " rowbuf[ cr + 3 ] ), /r/n" " ( REALTYPE4 ) ( in2[c * COLSROWS + cr ], /r/n" " in2[c * COLSROWS + cr + 1 ], /r/n" " in2[c * COLSROWS + cr + 2 ], /r/n" " in2[c * COLSROWS + cr + 3 ] ) ); /r/n" " out[ r * COLS2 + c ] = sum; /r/n" " } /r/n" "} /r/n" ;
清单8。使用float4内核的部分矢量化（仅限内部循环）
完整源代码的文件为 matr_mul_vect.mq5。当然，COLSROWS 参数应可被 4 整除。
CPU：
2012.05.27 21:28:16 matr_mul_vect (EURUSD,H1) CPUTime / GPUTotalTime = 18.657 2012.05.27 21:28:16 matr_mul_vect (EURUSD,H1) OpenCL total: time = 4.945 sec. 2012.05.27 21:28:16 matr_mul_vect (EURUSD,H1) read = 4000000 elements 2012.05.27 21:28:11 matr_mul_vect (EURUSD,H1) CPUTime = 92.259 2012.05.27 21:26:38 matr_mul_vect (EURUSD,H1) 1st OCL martices mul: device = 0; ROWS1 = 2000; COLSROWS = 2000; COLS2 = 2000 2012.05.27 21:26:38 matr_mul_vect (EURUSD,H1) =======================================
GPU：
2012.05.27 21:21:30 matr_mul_vect (EURUSD,H1) CPUTime / GPUTotalTime = 78.079 2012.05.27 21:21:30 matr_mul_vect (EURUSD,H1) OpenCL total: time = 1.186 sec. 2012.05.27 21:21:30 matr_mul_vect (EURUSD,H1) read = 4000000 elements 2012.05.27 21:21:28 matr_mul_vect (EURUSD,H1) CPUTime = 92.602 2012.05.27 21:19:55 matr_mul_vect (EURUSD,H1) 1st OCL martices mul: device = 1; ROWS1 = 2000; COLSROWS = 2000; COLS2 = 2000 2012.05.27 21:19:55 matr_mul_vect (EURUSD,H1) =======================================
令人惊讶的是，即使是最初的矢量化，在GPU上也取得了良好的效果；虽然不是太大，但是增益似乎仍然是10%左右。
在内核中继续矢量化：将“昂贵的”raltype4矢量类型转换操作连同显式矢量组件的规范转换为外部辅助循环，填充专用变量rowbuf[]。内核仍然没有变化。
const string clSrc = "#define COLS2 " + i2s( COLS2 ) + " /r/n" "#define COLSROWS " + i2s( COLSROWS ) + " /r/n" "#define REALTYPE float /r/n" "#define REALTYPE4 float4 /r/n" " /r/n" "__kernel void matricesMul( __global REALTYPE *in1, /r/n" " __global REALTYPE *in2, /r/n" " __global REALTYPE *out ) /r/n" "{ /r/n" " int r = get_global_id( 0 ); /r/n" " REALTYPE4 rowbuf[ COLSROWS / 4 ]; /r/n" " for( int col = 0; col < COLSROWS / 4; col ++ ) /r/n" " { /r/n" " rowbuf[ col ] = ( REALTYPE4 ) ( in1[r * COLSROWS + 4 * col ], /r/n" " in1[r * COLSROWS + 4 * col + 1 ], /r/n" " in1[r * COLSROWS + 4 * col + 2 ], /r/n" " in1[r * COLSROWS + 4 * col + 3 ] ); /r/n" " } /r/n" " /r/n" " REALTYPE sum; /r/n" " /r/n" " for( int c = 0; c < COLS2; c ++ ) /r/n" " { /r/n" " sum = 0.0; /r/n" " for( int cr = 0; cr < COLSROWS / 4; cr ++ ) /r/n" " sum += dot( rowbuf[ cr ], /r/n" " ( REALTYPE4 ) ( in2[c * COLSROWS + 4 * cr ], /r/n" " in2[c * COLSROWS + 4 * cr + 1 ], /r/n" " in2[c * COLSROWS + 4 * cr + 2 ], /r/n" " in2[c * COLSROWS + 4 * cr + 3 ] ) ); /r/n" " out[ r * COLS2 + c ] = sum; /r/n" " } /r/n" "} /r/n" ;
清单9。摆脱主内核循环中的“昂贵”类型转换操作
注：内部（和辅助）循环计数器的最大计数值降低了四倍，现在第一个数组所需的读取操作比以前少了四倍——显然，读取已成为矢量操作。
CPU：
2012.05.27 22:41:43 matr_mul_vect_v2 (EURUSD,H1) CPUTime / GPUTotalTime = 24.480 2012.05.27 22:41:43 matr_mul_vect_v2 (EURUSD,H1) OpenCL total: time = 3.791 sec. 2012.05.27 22:41:43 matr_mul_vect_v2 (EURUSD,H1) read = 4000000 elements 2012.05.27 22:41:39 matr_mul_vect_v2 (EURUSD,H1) CPUTime = 92.805 2012.05.27 22:40:06 matr_mul_vect_v2 (EURUSD,H1) 1st OCL martices mul: device = 0; ROWS1 = 2000; COLSROWS = 2000; COLS2 = 2000 2012.05.27 22:40:06 matr_mul_vect_v2 (EURUSD,H1) =======================================
GPU：
2012.05.27 22:35:28 matr_mul_vect_v2 (EURUSD,H1) CPUTime / GPUTotalTime = 185.605 2012.05.27 22:35:28 matr_mul_vect_v2 (EURUSD,H1) OpenCL total: time = 0.499 sec. 2012.05.27 22:35:28 matr_mul_vect_v2 (EURUSD,H1) read = 4000000 elements 2012.05.27 22:35:27 matr_mul_vect_v2 (EURUSD,H1) CPUTime = 92.617 2012.05.27 22:33:54 matr_mul_vect_v2 (EURUSD,H1) 1st OCL martices mul: device = 1; ROWS1 = 2000; COLSROWS = 2000; COLS2 = 2000 2012.05.27 22:33:54 matr_mul_vect_v2 (EURUSD,H1) =======================================
算术吞吐量：
throughput_arithmetic_CPU_OCL = 16 000000000 / 3.791 ~ 4.221 GFlops. throughput_arithmetic_GPU_OCL = 16 000000000 / 0.499 ~ 32.064 GFlops.
可以看出，CPU性能发生了很大的变化，对于GPU来说，几乎是革命性的变化。有关源代码，请参阅matr_mul_vect_v2.mq5。
我们只使用宽度为8的向量来执行与内核变量相关的相同操作。作者的决定可以用256位的GPU内存带宽（即32字节和8个浮点数字）来解释；因此，同时处理8个浮点数据似乎很自然（相当于并行使用float8）。
记住：在这种情况下，colsrow值应该可以被8整除。这是正常的要求，因为更精细的优化将需要更具体的数据设置。
const string clSrc = "#define COLS2 " + i2s( COLS2 ) + " /r/n" "#define COLSROWS " + i2s( COLSROWS ) + " /r/n" "#define REALTYPE float /r/n" "#define REALTYPE4 float4 /r/n" "#define REALTYPE8 float8 /r/n" " /r/n" "inline REALTYPE dot8( REALTYPE8 a, REALTYPE8 b ) /r/n" "{ /r/n" " REALTYPE8 c = a * b; /r/n" " REALTYPE4 _1 = ( REALTYPE4 ) 1.; /r/n" " return( dot( c.lo + c.hi, _1 ) ); /r/n" "} /r/n" " /r/n" "__kernel void matricesMul( __global REALTYPE *in1, /r/n" " __global REALTYPE *in2, /r/n" " __global REALTYPE *out ) /r/n" "{ /r/n" " int r = get_global_id( 0 ); /r/n" " REALTYPE8 rowbuf[ COLSROWS / 8 ]; /r/n" " for( int col = 0; col < COLSROWS / 8; col ++ ) /r/n" " { /r/n" " rowbuf[ col ] = ( REALTYPE8 ) ( in1[r * COLSROWS + 8 * col ], /r/n" " in1[r * COLSROWS + 8 * col + 1 ], /r/n" " in1[r * COLSROWS + 8 * col + 2 ], /r/n" " in1[r * COLSROWS + 8 * col + 3 ], /r/n" " in1[r * COLSROWS + 8 * col + 4 ], /r/n" " in1[r * COLSROWS + 8 * col + 5 ], /r/n" " in1[r * COLSROWS + 8 * col + 6 ], /r/n" " in1[r * COLSROWS + 8 * col + 7 ] ); /r/n" " } /r/n" " /r/n" " REALTYPE sum; /r/n" " /r/n" " for( int c = 0; c < COLS2; c ++ ) /r/n" " { /r/n" " sum = 0.0; /r/n" " for( int cr = 0; cr < COLSROWS / 8; cr ++ ) /r/n" " sum += dot8( rowbuf[ cr ], /r/n" " ( REALTYPE8 ) ( in2[c * COLSROWS + 8 * cr ], /r/n" " in2[c * COLSROWS + 8 * cr + 1 ], /r/n" " in2[c * COLSROWS + 8 * cr + 2 ], /r/n" " in2[c * COLSROWS + 8 * cr + 3 ], /r/n" " in2[c * COLSROWS + 8 * cr + 4 ], /r/n" " in2[c * COLSROWS + 8 * cr + 5 ], /r/n" " in2[c * COLSROWS + 8 * cr + 6 ], /r/n" " in2[c * COLSROWS + 8 * cr + 7 ] ) ); /r/n" " out[ r * COLS2 + c ] = sum; /r/n" " } /r/n" "} /r/n" ;
清单10。使用宽度为8的向量进行核矢量化
我们必须在内核代码中插入一个内联函数dot8（），它允许计算8个向量的无向积。在opencl中，标准函数dot（）只能计算宽度不超过4的向量的无向积。有关源代码，请参阅matr_mul_vect_v3.mq5。
CPU：
2012.05.27 23:11:47 matr_mul_vect_v3 (EURUSD,H1) CPUTime / GPUTotalTime = 45.226 2012.05.27 23:11:47 matr_mul_vect_v3 (EURUSD,H1) OpenCL total: time = 2.200 sec. 2012.05.27 23:11:47 matr_mul_vect_v3 (EURUSD,H1) read = 4000000 elements 2012.05.27 23:11:45 matr_mul_vect_v3 (EURUSD,H1) CPUTime = 99.497 2012.05.27 23:10:05 matr_mul_vect_v3 (EURUSD,H1) 1st OCL martices mul: device = 0; ROWS1 = 2000; COLSROWS = 2000; COLS2 = 2000 2012.05.27 23:10:05 matr_mul_vect_v3 (EURUSD,H1) =======================================
GPU：
2012.05.27 23:20:05 matr_mul_vect_v3 (EURUSD,H1) CPUTime / GPUTotalTime = 170.115 2012.05.27 23:20:05 matr_mul_vect_v3 (EURUSD,H1) OpenCL total: time = 0.546 sec. 2012.05.27 23:20:05 matr_mul_vect_v3 (EURUSD,H1) read = 4000000 elements 2012.05.27 23:20:04 matr_mul_vect_v3 (EURUSD,H1) CPUTime = 92.883 2012.05.27 23:18:31 matr_mul_vect_v3 (EURUSD,H1) 1st OCL martices mul: device = 1; ROWS1 = 2000; COLSROWS = 2000; COLS2 = 2000 2012.05.27 23:18:31 matr_mul_vect_v3 (EURUSD,H1) =======================================
结果是出乎意料的：CPU上的运行时间几乎是以前的两倍，而GPU上的运行时间稍长一些，尽管float8有足够的总线宽度（相当于256位）用于hd 4870。这里，我们再次需要转到GPU Caps Viewer。
有关说明，请参见图18:
中参数列表中的倒数第二行。
cl_设备_首选_向量_宽度_浮点：4
参见
中的OpenCL规范，您将在第37页表4.3的最后一列中看到关于此参数的以下文本：
可放入向量的内置标量类型的首选原始向量宽度大小。矢量宽度定义为可以存储在矢量中的缩放元素的数量。
因此，对于HD 4870，矢量floatn的首选矢量宽度是float4而不是float8。
让我们在这里结束内核优化循环。我们可以对此进行更多的讨论，但它仅限于本文的篇幅，不会深入讨论。
总结
本文描述了一些优化功能，但至少您需要了解一些内核用于执行以启动这些功能的基本硬件。
获得的数据远未达到最大值，但即便如此，仍建议充分利用现有资源（终端开发人员实现的Opencl API不允许控制某些对优化非常重要的参数，尤其是工作组的大小）。主机程序执行带来的收益是相当可观的：GPU执行和CPU（尽管没有完全优化）带来的收益。顺序程序增益比约为200:1。
我非常感谢元驱动程序的宝贵建议和在我有一个独立的GPU之前利用它的机会。
附件目录：
- 母女。mq5—主机上的初始序列程序，采用二维数据形式；
- matr_mul_1dim.mq5——第一个具有线性数据格式并绑定到mql5 opencl api的内核实现；
- Matr_Mul_1dim_合并-与包含全局内存访问的内核结合；
- matr_mul_sum_local-用于计算无向积的专用变量，不再访问存储在全局内存输出数组中的计算单元；
- matr_mul_row_calc—计算内核中的整个输出矩阵行；
- matr-mul-row-in-private-rows迁移到第一组私有内存；
- matr_mul_col_local.mq5—迁移到第二个本地内存阵列的列；
- matr_mul_vect.mq5——内核的第一个矢量化（使用float4，只在主循环中使用内部子循环）；
- matr_mul_vect_v2.mq5—摆脱主循环中的“昂贵”数据转换操作；
- matr_mul_vect_v3.mq5-使用宽度为8的向量执行矢量化。
本文由MetaQuotes Software Corp.翻译自俄语原文
，网址为https://www.mql5.com/ru/articles/407。
MyFxtop迈投(www.myfxtop.com)-靠谱的外汇跟单社区，免费跟随高手做交易！
免责声明：本文系转载自网络，如有侵犯，请联系我们立即删除，另：本文仅代表作者个人观点，与迈投财经（www.myfxtop.cn）无关。其原创性以及文中陈述文字和内容未经本站证实，对本文以及其中全部或者部分内容、文字的真实性、完整性、及时性本站不作任何保证或承诺，请读者仅作参考，并请自行核实相关内容。
著作权归作者所有。
商业转载请联系作者获得授权，非商业转载请注明出处。
