OpenCL: 실제 문제 - 페이지 8

[Roffild]

Mathemat :

자, 자신의 코드를 보십시오. 그리고 마지막 줄에서 240을 18로 나눕니다(이것은 지도의 단위입니다).

당신은 분명히 뭔가를 혼동하고 있습니다. 논란이 되는 부분은 다음과 같다.

   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]);

출력: 전역=30 지역=1

그리고 버퍼를 생성할 때 240바이트.

[Sceptic Philozoff]

Roffild :

당신은 분명히 뭔가를 혼동하고 있습니다. 논란이 되는 부분은 다음과 같다.

출력: 전역=30 지역=1

그리고 버퍼를 생성할 때 240바이트.

전역 = 240. 인쇄

global_work_size[ 0 ]

그리고 local_work_size[0] = (단위) 240/18 = 13

PS 네, 맞습니다. 죄송합니다. 조금 혼란스럽습니다.

local_work_size[0] = (단위) 30/18 = 1 단위=28.

[Sceptic Philozoff]

다시 한 번, 로 필드 :

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

병렬 작업(Ocl에 대한 첫 번째 분기의 MetaDriver 스크립트 참조)을 통해 GPU에서 최대 1000개 이상의 속도 향상을 달성할 수 있습니다(MQL5의 CPU에서 1개 스레드에서 실행하는 것과 비교). 찾을 수 없으면 버릴 수 있습니다. 지도에서 테스트해 보세요.

[Roffild]

버퍼와 속도를 알아냈습니까?

그리고 UNITS 등에 대해 AMD CodeXL을 통해 파악하는 것이 좋습니다. 성능 그래프가 아름답습니다.

AMD CodeXL 자체는 버그가 있는 감염이지만 이것이 없으면 결론을 내리기가 어렵습니다.

테스터가 CPU를 허용하거나 버퍼 수 * 0.353ms보다 더 오래 걸리는 작업이 나타날 때까지 OpenCL을 더 이상 다루지 않을 것입니다.

추신

마침내 코드 최적화를 완료했으며 이제 최종 버전은 33초(최적화 전 320초, "OpenCL 스타일"에서 55초) 만에 테스트를 통과합니다.

[Sceptic Philozoff]

Roffild : ?

그리고 무언가를 이해하기 위해 거기에 무엇이 있습니다. 이것은 느린 작업임이 분명합니다. 결론은 커널 내부의 작업을 늘리는 것입니다(코드에 작업량이 너무 적음).

그리고 더 현대적인 vidyahi를 가지고, 그들은 이것으로 더 좋아진 것 같습니다.

AMD CodeXL 자체는 버그가 있는 감염이지만 이것이 없으면 결론을 내리기가 어렵습니다.

Intel 유틸리티도 매우 유용하지만 Intel 스톤에는 유용합니다. 글쎄, 커널에서 가장 명백한 오류를 잡기 위해.

추신: 마침내 코드 최적화를 완료했으며 이제 최종 버전은 33초(최적화 전 320초, "OpenCL 스타일"에서 55초) 만에 테스트를 통과합니다.

이미 훨씬 낫습니다.

[Roffild]

오늘은 숫자가 1비트인 배열을 생성해야 했습니다.

동시에 OpenCL에서 연습했습니다.

global_work_size 및 local_work_size를 계산하는 흥미로운 방법의 데모로 코드를 게시하고 있습니다. 아이디어 자체는 IntrotoOpenCL.pdf (나는 사본이 있다)에서 가져왔지만 수정했다.

 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);
}