OpenCL (4 стр)

Сейчас любуюсь oclNbody (проблема n тел), у меня на девайсе:

Device GeForce 8800 GT:
  CL_DEVICE_TYPE:      CL_DEVICE_TYPE_GPU
  CL_DEVICE_MAX_COMPUTE_UNITS:    14
  CL_DEVICE_MAX_WORK_ITEM_SIZES:  512 / 512 / 64
  CL_DEVICE_MAX_WORK_GROUP_SIZE:  512
  CL_DEVICE_MAX_CLOCK_FREQUENCY:  1 MHz
  CL_DEVICE_IMAGE_SUPPORT:    1
  CL_DEVICE_GLOBAL_MEM_SIZE:    1023 MByte
  CL_DEVICE_LOCAL_MEM_SIZE:    16 KByte

он дал такие показатели:

OpenCL for GPU Nbody Demo (8192 bodies): 75 fps | 5.0 BIPS | 100.7 GFLOP/s

О. Федор
Круто :)
Да, 14 ядер по 8 скаляров в каждом, всё правильно.

Не, теоретически должно быть 300-400 GFLOP/s. Но 100 тоже хорошо.

О. Федор
Это для MAD.
Делим на 2, получаем загрузку около 50%.

А где для Вин7 дрова брать? :(

Ура, теперь компилирую в vs 2003 !
Просто в .vcproj сменить цифирьку с 9.00 на 7.10 оказалось недостаточным, нужно было файл подчистить.

Но сейчас проверил как пройдет при вызове из кернела вот такой простенький код

template <typename T, int n>
void test(void)
{}

elements of input arrays
__kernel void uniformUpdate(
    __global uint4 *d_Data,
    __constant uint *c_Scans
){
  test<float, 4>();      //вот такая шаблонная функция
    uint4 data4 = d_Data[get_global_id(0)];
    data4[0] += c_Scans[get_group_id(0)];
    data4[1] += c_Scans[get_group_id(0)];
    data4[2] += c_Scans[get_group_id(0)];
    data4[3] += c_Scans[get_group_id(0)];
    d_Data[get_global_id(0)] = data4;
}

Получилось весьма печально: компилируется, но при исполнении сообщает Error #11, не может произвести Building program.
Похоже, что ребята из Кроноса слишком серьезно отнеслись к выполнению своего обещания следовать стандарту C99. В OpenCL будем иметь для кернелей чистый C, хотя на хосте в полный рост используется C++.
Поковыряюсь еще малость и вернусь к CUDA. Там нет таких ограничений и можно и для кернелей очень многое использовать из  C++.

Executor
> А где для Вин7 дрова брать? :(

Для Висты CUDA драйвер вроде есть. Загрузи один из этих файлов

19.06.2009  22:39        52 949 377 nvcomputesdk_opencl1.0_conformant_winxp_32.exe
http://www.rapidshare.ru/1076969

19.06.2009  22:39        54 003 751 nvcomputesdk_opencl1.0_conformant_winxp_64.exe
http://www.rapidshare.ru/1076974

Думаю будет работать.

Сейчас посмотрел. Для Висты точно есть поддержка. В каталоге

...\NVIDIA GPU Computing SDK\OpenCL\common\lib\WinVista32

есть файлы

libOpenCL32.dll
libOpenCL32.lib

Этого видимо достаточно.

О. Федор
> ...\NVIDIA GPU Computing SDK\OpenCL\common\lib\WinVista32
а если у мну Vista x64 ? Вопщем походу надо ждать офф. релиза от nVidia :(

Ждём поддержки от ATI. Но что-то у них на сайте глухо об этом.

Booster
> Ждём поддержки от ATI. Но что-то у них на сайте глухо об этом.
Да уж... Не до того им сейчас, им DirectX 11 проталкивать голова болит ;)

0r@ngE
> а если у мну Vista x64 ? Вопщем походу надо ждать офф. релиза от nVidia :(

Ох, я не знаю. Тебе нужен просто кудавский драйвер. Мне кажется на сайте нвидия я его видел.

Booster
> Ждём поддержки от ATI. Но что-то у них на сайте глухо об этом.

Я тоже, как узнал о OpenCL тоже стал даже подумывать о новом ноуте с ATI картой, тем более, что как я читал  топовые радеоны имеют сейчас до 1600 процессоров. Но когда опробовал OpenCL, сразу отбросил эти мысли. Без шаблонов и возможности симулирования рекурсивных вызовов новыми инстанциями шаблонов, программирование GPU очень сложно, поскольку рекурсивные по сути алгоритмы придется приводить к итерационному виду.

0r@ngE
> а если у мну Vista x64 ? Вопщем походу надо ждать офф. релиза от nVidia :(

Кстати, компилятор OpenCL для Vista x64 есть. Это точно.

О. Федор
> Без шаблонов и возможности симулирования рекурсивных вызовов новыми инстанциями
> шаблонов, программирование GPU очень сложно, поскольку рекурсивные по сути
> алгоритмы придется приводить к итерационному виду.
А почему такая необходимость в этом? По-моему рекурсивные алгоритмы не так часто встречаются даже при кодировании CPU.

0r@ngE
> Не до того им сейчас, им DirectX 11 проталкивать голова болит ;)
Вроде они уже и не собираются проталкивать. Вообще в чём плюс вычислительных шейдеров перед этим же OpenCL?

Booster
> А почему такая необходимость в этом? По-моему рекурсивные алгоритмы не так
> часто встречаются даже при кодировании CPU.

Не только рекурсии. Тонкая настройка программы тоже.
Недавно поднял производительнность в одной программке на 40% только изменив параметр шаблона. В OpenCL для этого пришлось бы использовать дефайн.
В общем, хочу шаблоны!