OpenCL：真正的挑战 - 页 5

[Ilyas]

向SD 创建一个请求 并附上代码（你可以通过PM来做），我将分析你的代码。

[Roffild]

你到底对我的哪部分代码感兴趣？我有很多对不同文件的依赖性。

我现在遇到的问题是在测试器的1次勾选中只能写入和读取缓冲区，对于检查来说已经足够了。

#property copyright ""
#property link      ""

int hcontext, hprogram, hkernel, hbuffer[5];

void InitGlobal()
{
   for (int cdev = (int)CLGetInfoInteger(0, CL_DEVICE_COUNT)-1; cdev>-1; cdev--)
   {
      string name;
      CLGetInfoString(cdev, CL_DEVICE_NAME, name);
      Print("Device #",cdev," = ",name);
   }
   
   string source =
"kernel void tester(global double *price, global double *result)                                   \r\n"
"{                                                                                                 \r\n"
"   int global_index = get_global_id(0);                                                           \r\n"
"   result[global_index] = price[global_index] / global_index;                                     \r\n"
"}                                                                                                 \r\n"
;
   
   hcontext = CLContextCreate(CL_USE_GPU_ONLY);
   string build_log;
   hprogram = CLProgramCreate(hcontext, source, build_log);
   Print("build log = ", build_log);
   hkernel = CLKernelCreate(hprogram, "tester");
}

void DeinitGlobal()
{
   CLBufferFree(hbuffer[0]);
   CLBufferFree(hbuffer[1]);
   
   CLKernelFree(hkernel);
   CLProgramFree(hprogram);
   CLContextFree(hcontext);
}

int OnInit()
{
   InitGlobal();
   return(0);
}

void OnDeinit(const int reason)
{
   DeinitGlobal();
}

// Скрипт, в отличии от эксперта, можно дебажить на выходных :)
//void OnStart() {  InitGlobal();
void OnTick() {
   double price[30];
   CopyClose(_Symbol,_Period,0,ArraySize(price),price);
   
   static bool firststart = true;
   if (firststart)
   {
      firststart = false;
      uint bufsize = sizeof(price);
      Print("Размер буфера в байтах =",bufsize);
      hbuffer[0] = CLBufferCreate(hcontext, bufsize, CL_MEM_READ_ONLY);
      hbuffer[1] = CLBufferCreate(hcontext, bufsize, CL_MEM_WRITE_ONLY);
      
      CLSetKernelArgMem(hkernel, 0, hbuffer[0]);
      CLSetKernelArgMem(hkernel, 1, hbuffer[1]);
   }
   
   // А вот здесь не хватает clGetMemObjectInfo(buffer, CL_MEM_SIZE) для проверки размера.
   
   CLBufferWrite(hbuffer[0], price);
   
   uint units = (uint)CLGetInfoInteger(hcontext, CL_DEVICE_MAX_COMPUTE_UNITS);
   uint global_work_offset[] = {0};
   uint global_work_size[1];
   uint local_work_size[1];
   global_work_size[0] = ArraySize(price);
   local_work_size[0] = global_work_size[0] / units;
   bool exec = CLExecute(hkernel, 1, global_work_offset, global_work_size, local_work_size); // async
   if (exec == false) Print("Error in ",__FUNCSIG__," CLExecute: ",GetLastError());
   
   CLBufferRead(hbuffer[1], price);
   
   if (MQL5InfoInteger(MQL5_PROGRAM_TYPE) == PROGRAM_SCRIPT) DeinitGlobal();
}

按脚本运行。

2013.10.30 18:55:40 OpenCL_buffer_test (EURUSD,H1) 设备 #1 = AMD Phenom(tm) II X4 925 处理器
2013.10.30 18:55:40 OpenCL_buffer_test (EURUSD,H1) 设备#0 = Cypress
2013.10.30 18:55:40 OpenCL_buffer_test (EURUSD,H1) OpenCL: GPU设备 "Cypress "被选中
2013.10.30 18:55:40 OpenCL_buffer_test (EURUSD,H1) build log =
2013.10.30 18:55:40 OpenCL_buffer_test (EURUSD,H1) 缓冲区大小，字节数=240

从2013.01.09到2013.10.10在M5上运行专家，使用 "OHLC on M1"。

2013.10.30 19:01:44 核心 1 EURUSD,M5: 测试 experts\OpenCL_buffer_test.ex5 从 2013.01.09 00:00 到 2013.10.10 00:00 开始
2013.10.30 19:01:44 核心 1 2013.01.09 00:00:00 设备 #0 = Cypress
2013.10.30 19:01:44 Core 1 2013.01.09 00:00:00 OpenCL: GPU设备'Cypress'被选中
2013.10.30 19:01:44 核心 1 2013.01.09 00:00:00 建立日志 =
2013.10.30 19:01:44 核心 1 2013.01.09 00:00:00 缓冲区大小（字节数）=240
2013.10.30 19:04:55 Core 1 EURUSD,M5: 1108637 ticks (55953 bars) generated within 192443 ms (total bars in history 131439, total time 192521 ms)
2013.10.30 19:04:55 核心 1 294 Mb 已用内存

请注意，测试器中只有1个设备。

如果

   //CLBufferRead(hbuffer[1], price);

然后

2013.10.30 19:16:00 核心 1 EURUSD,M5: 在88218毫秒内产生了1108637点（55953条）（历史总条数131439，总时间88297毫秒）。

[Ilyas]

读取是需要的；事实上，它是等待内核终止的函数，因为CLExecute 只排队执行作业，然后将控制权返回给MQL程序，而不等待它终止。

[Ilyas]

具体到这个例子中，使用OpenCL的 优势被缓冲区复制的开销所吞噬。

如果有必要对OHLC数据进行计算，那么就必须进行节约性写入，即提前创建一个较大的缓冲区，只有在新数据到达时才覆盖这些新数据，并告诉内核缓冲区的新起点和大小。

[Roffild]

即使我们设法优化了OHLC的传输（我们将使用CLSetKernelArg 来传输最后一个条形图），我们在读取结果缓冲区时仍然会崩溃。

2013.10.31 19:24:13 核心 1 EURUSD,M5: 1108637点（55953条）在114489毫秒内产生（历史总条数131439，总时间114598毫秒）。

(将有CLBufferWrite(hbuffer[0], price);的一行移到IF下面)

[Sceptic Philozoff]

Roffild: 呃...关于在GPU上使用OpenCL提高速度的文章，结果是一个童话故事，因为它们没有真正解决手头的任务。

那么，谁能阻止你这样做呢？去写一些真实的东西吧，这不会是一个童话故事。但要尝试找到一个例子，这样加速就会发生。这是最困难的部分。

如果你在谈论我的文章...好吧，我在写一本入门书。而矩阵乘法是一个相当有用的操作。

P.S. 顺便说一下，如果你的CPU是英特尔的，在它 上面模拟x86内核要比在竞争对手的CPU上模拟快得多。这是在你重新计算每个核心的情况下。

HD5850：基本上是一个相当体面的卡，但现代卡更好--不仅是由于更多的苍蝇，而且还由于OpenCL优化。例如，全局内存访问时间大大减少。

P.P.S. OpenCL并不是万能的，它是一个可行的工具，在一些罕见的情况下可以显著提高速度。而在其他不那么方便的情况下，加速效果也不那么令人印象深刻--如果有的话。

[Vladimir Gomonov]

Roffild:
呃...关于在GPU上使用OpenCL提高速度的文章，结果是一个童话，因为他们并没有真正处理真正的任务。

不是这样的。 童话故事是 "任何算法都可以在OpenCL中加速"。 不是任何算法。

关于OpenCL的第一个主题甚至很好地描述了一个算法必须具备的标准，以便具有ocl加速潜力。

祝你好运。

[Roffild]

这个说法不是关于计算速度的--有2倍的速度提升（0.02毫秒vs 0.05毫秒）。

声称文章中没有任何信息。

• 即使是一个小的缓冲区的读/写延迟=0.35300毫秒--这就是大多数算法转换为OpenCL的无效之处!
• 测试员没有为OpenCL选择CPU - 这一点在任何地方都没有报告!

我可能是第一个想以牺牲GPU为代价来加快测试速度的人，在看过承诺后...

MetaDriver: OpenCL上的第一个主题甚至很好地描述了一个算法必须具备的标准，即具有ocl加速潜力。

再读一下我的帖子。

主要标准： 在 "OpenCL-style "中执行MQL代码的时间应该超过1个tick的时间=Number of_Buffers * 0.35300 ms。

为了找出MQL中算法的速度，其精确度为微秒（1000微秒=1毫秒），你必须在测试器和Total_Time / Number_of_Ticks（我的顶帖）中运行数次。

如果没有缓冲区延迟，我的代码将在30秒内通过测试 - 这比 "OpenCL风格 "的MQL（55秒）快2倍，比普通代码（320秒）快11倍。

还有什么其他标准？

[Sceptic Philozoff]

Roffild: 声称文章中没有任何信息。

• 即使是一个小的缓冲区的读/写延迟=0.35300毫秒--这就是大多数算法转换为OpenCL的意义所在!

从你与OpenCL打交道的经验来看，你一定已经明白，不是每一种算法都能直接加速。这里的主要问题之一是尽量减少全局内存访问。

顺便说一句，我现在要解决一个类似的问题，即随机访问全局设备内存的问题（太私人的这种随机访问，而且是他妈的开销）。一旦我的大脑恢复正常，我就会解决这个问题。

测试员没有为OpenCL选择CPU - 这一点在任何地方都没有报告!

写信给服务台，证明对这种功能的需要。

如果不使用测试器，就已经完成了（这是我的应用）。而且我还没有检查过测试器。

[Roffild]

Mathemat:

前面已经写过，不是每个算法都能直接加速。你必须在这里动脑筋，主要问题之一是尽量减少全局内存的访问。

好了，现在我必须解决一个类似的随机访问全局内存的问题（这个随机访问太频繁了）。一旦我的大脑开始工作，我就会解决这个问题。

是时候动动脑子了，因为0.35300毫秒正是指clEnqueue[Read/Write]Buffer()，而不是指内核 内的全局内存访问。

第二个问题可以通过优化内核本身来解决，而第一个问题是一个铁的限制。