Комментировать

GTX480 новые результаты

не мог заснуть, по-этому решил сделать тест на bandwidth dtod:
код ядер:

  1. #pragma OPENCL EXTENSION cl_khr_fp64 : enable
  2. #define AEQB(FUNCNAME,VECTORTYPE) \
  3. __kernel void FUNCNAME(__global VECTORTYPE *vA,__global VECTORTYPE *vB,const unsigned int n) \
  4. { \
  5.     uint tid = get_global_id(0); \
  6.     if(tid<n) \
  7.     { \
  8.         vA[tid]=vB[tid]; \
  9.     } \
  10. }
  11. AEQB(DAEqB,double)
  12. AEQB(SAEqB,float)

Размер группы 512 (определилось недолгим подбором, много времени строить графики нет..)
С увеличением N, увеличивалась и скорость передачи. Но сильно большое N не получиться взять, так как OpenCL вроде ограничивает размер одного буфера.

Время выполнения замерялось OpenCL'скими средствами.

  1. clGetEventProfilingInfo(events[0],CL_PROFILING_COMMAND_START,sizeof(cl_ulong),&startTime,0);
  2. clGetEventProfilingInfo(events[0],CL_PROFILING_COMMAND_END,sizeof(cl_ulong),&endTime,0);

У меня есть предположение, что в это время входит и коммуникация с хостом - передача ядра, параметров запуска и т.п., по-этому с увеличением N, увеличивался показатель bandwidth измеряемый мной.
Вот результаты

SAEqB N=134217728 : 124.4 GB/s
DAEqB N=67108864 : 136.9 GB/s

Наверное увеличивая размер передаваемых данных можно получить больший показатель, но заниматься этим времени нет (если просто добавить цикл в ядро, то мы уже будем попадать в кэш, надо учитывать это при составлении теста - желательно ещё смотреть что говорит профилировщик).

Сделал грубый замер - запустил ядро с N=1 - время С=0.000015с (NVidia выдаёт в CL_DEVICE_PROFILING_TIMER_RESOLUTION разрешение 1000ns),
запускал ядра с меньшим N:
SAEqB N=2097152
DAEqB N=1048576
показатель в нескольких вариантах - с вычетом C и без.

Вычет C Raw
SAEqB 125.5GB/s 112.1GB/s
DAEqB 137.9GB/s 122.2GB/s

Итого, результат с вычетом C достаточно близкий к запуску при большом N(без вычета).

Для перевода из байт в гигабайты делил на 2^30