OpenCL: desafios reais - página 8

[Roffild]

Mathemat:

Veja o seu próprio código: E depois, na última linha, você mesmo divide 240 por 18 (isto é, unidades para o seu cartão).

Está obviamente confuso com alguma coisa. Aqui está a peça controversa:

   double price[30];

   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); /// <-- здесь же НЕ sizeof(price) вообще-то!
   local_work_size[0] = global_work_size[0] / units;

Print("global=",global_work_size[0]," local=",local_work_size[0]);

Conclusão: global=30 local=1

E 240 bytes é exactamente quando se cria o tampão.

[Sceptic Philozoff]

Roffild:

Está obviamente confuso com alguma coisa. Aqui está uma peça controversa:

Produção: global=30 local=1

E 240 bytes exactamente quando se cria o tampão.

global = 240. Imprimir

global_work_size[0]

E local_work_size[0] = (uint) 240/18 = 13

P.S. Sim, acertou. Perdão. Fiquei um pouco confuso.

local_work_size[0] = (uint) 30/18 = 1. E eu tenho o mesmo, desde unidades=28.

[Sceptic Philozoff]

Mais uma vez, Roffild:

Mathemat: Давай тупо прикинем. 18 задач, выполняемых одновременно на мухах GPU, - это максимум то, что можно сделать на 4-5 нитках CPU. А CPU на x86 эмуляции может организовать гораздо больше ниток. Во всяком случае, если это Intel. Мой бывший Pentium G840 (2 ядра) дал ускорение примерно в 70 раз - на двух unit'ах! Я уже не говорю о том, что вытворяет мой текущий... условно говоря, i7.

Uma tarefa bem paralela (ver os scripts do MetaDriver da primeira thread ocl) pode atingir velocidades até 1000 ou mais na GPU (em comparação com a execução de 1 thread na CPU no MQL5). Se não o conseguir encontrar - posso enviar-lho, pode testá-lo no seu cartão.

[Roffild]

Já descobriu o tampão e a sua velocidade?

É melhor usar o AMD CodeXL para descobrir UNIDADES, etc. - tem bons gráficos de desempenho.

O próprio AMD CodeXL é defeituoso, mas é difícil tirar quaisquer conclusões sem ele.

Não vou usar OpenCL até que o testador me permita usar CPU ou até que execute uma tarefa que dure mais do que o Número de_buffers * 0,353 msec.

P.S.

Finalmente terminei a optimização do meu código e a variante final passa o teste em 33 segundos (320 segundos - antes da optimização, 55 segundos - "OpenCL-style").

[Sceptic Philozoff]

Roffild: Já descobriu o tampão e a sua velocidade?

Não há nada a descobrir. É evidente que se trata de uma operação lenta. Conclusão - aumentar o trabalho dentro do núcleo (há muito pouco dele no seu código).

E comprar uma placa de vídeo mais moderna, parece ter-se tornado melhor com ela.

O próprio AMD CodeXL é defeituoso, mas é difícil tirar quaisquer conclusões sem ele.

A utilidade da Intel é bastante útil também, mas para as pedras Intel. Bem, e por apanhar os erros mais óbvios no núcleo.

P.S. Afinal, acabei de optimizar o meu código e a variante final passa agora o teste em 33 segundos (320 segundos - antes da optimização, 55 segundos - "estilo OpenCL").

Já está muito melhor.

[Roffild]

Hoje precisava de gerar uma matriz com 1 bit em números.

Ao mesmo tempo, pratiquei com OpenCL.

Estou a afixar o código como uma demonstração de um método interessante para calcular o_tamanho_trabalho_global e o_tamanho_trabalho_local. A ideia em si é retirada de IntrotoOpenCL.pdf (tenho uma cópia), mas eu afinei-a.

void OnStart()
{
   const string source =
      "kernel void differ(const int sizearray, const int bits, global uchar *result)        \r\n"
      "{                                                                                    \r\n"
      "   size_t global_index = get_global_id(0);                                           \r\n"
      "   if (global_index >= sizearray) return; // проверка границ, когда work>arraysize   \r\n"
      "   size_t xor = global_index;                                                        \r\n"
      "   uchar secc = 0;                                                                   \r\n"
      "   for (int bit = bits; bit>-1; bit--)                                               \r\n"
      "     if ((xor & ((size_t)1 << bit)) > 0) secc++;                                     \r\n"
      "   result[global_index] = secc;                                                      \r\n"
      "}                                                                                    \r\n"
   ;
   
   int hContext = CLContextCreate();
   string build_log = "";
   int hProgram = CLProgramCreate(hContext, source, build_log);
   Print("Error = ",build_log);
   int hKernel = CLKernelCreate(hProgram, "differ");
   
   uchar alldiff[1 << 17] = {0};
   CLSetKernelArg(hKernel, 0, ArraySize(alldiff));
   CLSetKernelArg(hKernel, 1, 17 /*bits*/);
   int hBuffer = CLBufferCreate(hContext, sizeof(alldiff), CL_MEM_WRITE_ONLY);
   CLSetKernelArgMem(hKernel, 2, hBuffer);
   CLBufferWrite(hBuffer, alldiff);
   
   /*uchar group_size[1024] = {0};
   uint deviceinfo_size = 0;
   CLGetDeviceInfo(hContext, CL_DEVICE_MAX_WORK_GROUP_SIZE, group_size, deviceinfo_size);
      for (int x = deviceinfo_size; x>=0; x--) Print(group_size[x]);
      Print("ch ",CharArrayToString(group_size));
   */ ///// CLGetDeviceInfo возвращает массив битов (шо за бред?)
   uint group_size = 256;
   
   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(alldiff);
   local_work_size[0] = global_work_size[0] / units;
   if (local_work_size[0] < 1) local_work_size[0] = 1;
   if (local_work_size[0] > group_size) local_work_size[0] = group_size;
   if (global_work_size[0] % local_work_size[0] != 0)
   {
      // увеличиваем global, чтобы global % local == 0
      // в самом kernel проверяется выход за границы
      global_work_size[0] = (int(global_work_size[0] / local_work_size[0]) +1) * local_work_size[0];
      // объяснение в
      // http://wiki.rac.manchester.ac.uk/community/OpenCL?action=AttachFile&amp;do=get&target=IntrotoOpenCL.pdf
   }
      Print("work=", global_work_size[0], " local=", local_work_size[0], " group=", group_size);
   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, alldiff);
   
   int hDump = FileOpen("alldiff.diff", FILE_ANSI|FILE_WRITE);
   for (int x = 0, xcount = ArraySize(alldiff); x < xcount; x++)
      FileWriteString(hDump, (string)alldiff[x]+",");
   FileClose(hDump);
   
   CLBufferFree(hBuffer);
   CLKernelFree(hKernel);
   CLProgramFree(hProgram);
   CLContextFree(hContext);
}