NVIDIA & CUDA

[Olej]

Так...
для этих игрищ нам нужен видео-девайс от NVIDIA:

И наиболее полную информацию о таком совместимом девайсе, при его наличии, даёт ещё один тест из того же SDK от NVIDIA:

Код: Выделить всё

[olej@nvidia release]$ ./deviceQuery
[deviceQuery] starting...
./deviceQuery Starting...

 CUDA Device Query (Runtime API) version (CUDART static linking)

Found 1 CUDA Capable device(s)

Device 0: "GeForce GT 520"
  CUDA Driver Version / Runtime Version          4.0 / 4.0
  CUDA Capability Major/Minor version number:    2.1
  Total amount of global memory:                 1023 MBytes (1072889856 bytes)
  ( 1) Multiprocessors x (48) CUDA Cores/MP:     48 CUDA Cores
  GPU Clock Speed:                               1.62 GHz
  Memory Clock rate:                             600.00 Mhz
  Memory Bus Width:                              64-bit
  L2 Cache Size:                                 65536 bytes
  Max Texture Dimension Size (x,y,z)             1D=(65536), 2D=(65536,65535), 3D=(2048,2048,2048)
  Max Layered Texture Size (dim) x layers        1D=(16384) x 2048, 2D=(16384,16384) x 2048
  Total amount of constant memory:               65536 bytes
  Total amount of shared memory per block:       49152 bytes
  Total number of registers available per block: 32768
  Warp size:                                     32
  Maximum number of threads per block:           1024
  Maximum sizes of each dimension of a block:    1024 x 1024 x 64
  Maximum sizes of each dimension of a grid:     65535 x 65535 x 65535
  Maximum memory pitch:                          2147483647 bytes
  Texture alignment:                             512 bytes
  Concurrent copy and execution:                 Yes with 1 copy engine(s)
  Run time limit on kernels:                     Yes
  Integrated GPU sharing Host Memory:            No
  Support host page-locked memory mapping:       Yes
  Concurrent kernel execution:                   Yes
  Alignment requirement for Surfaces:            Yes
  Device has ECC support enabled:                No
  Device is using TCC driver mode:               No
  Device supports Unified Addressing (UVA):      No
  Device PCI Bus ID / PCI location ID:           1 / 0
  Compute Mode:
     < Default (multiple host threads can use ::cudaSetDevice() with device simultaneously) >

deviceQuery, CUDA Driver = CUDART, CUDA Driver Version = 4.0, CUDA Runtime Version = 4.0, NumDevs = 1, Device = GeForce GT 520
[deviceQuery] test results...

- 48 float-point процессоров 1.62 GHz ... за ~$70

P.S. и таких подобных девайсов может стоять в системном блоке, как я пока разобрался из документации, может стоять несколько (3, 4), по числу доступных PCI слотов, что ещё позволит в разы распараллелить задачу... там, правда, с потреблением и теплом у этих девайсов тоже есть напряжёнка, и на 3 девайса ставят блок питания в ящик на 700W... однако, и зимой не замёрзнешь

[Olej]

О всём, что связано с параллельными вычислениями на GPU, буду писать сюда.

Если я правильно понимаю, то альтернативой CUDA-tools (единственной?) является OpenCL (Open Computing Language)?
Но это уже не только tools NVIDIA, но и Apple & AMD etc.
http://ru.wikipedia.org/wiki/OpenCL :

20 апреля 2009 года nVidia представила бета-драйвер и комплект для разработки программного обеспечения (SDK) с поддержкой открытого GPGPU-стандарта OpenCL. Данный бета-драйвер предназначен для разработчиков, участвующих в программе «OpenCL Early Access», которые уже с 20 апреля могут принять участие в тестировании бета-версии. Для участников программы «GPU Computing Registered Developers» бета-версия драйвера OpenCL будет доступна позже.[7][8][9]
5 апреля 2009 года компания AMD анонсировала доступность для загрузки beta-версии комплекта разработчика ATI Stream SDK v2.0 в который входит язык мультипроцессорного программирования OpenCL.
26 ноября 2009 года компания nVidia выпустила драйвер с поддержкой OpenCL 1.0 (rev 48).

От NVIDIA OpenCL, похоже, есть просто надстройкой над их CUDA... поскольку похожи они как близнецы братья...
http://www.nvidia.ru/object/cuda_opencl_new_ru.html :

Теперь вам доступны драйверы OpenCL v1.1 и примеры кода (июнь 2010)
Предварительные драйверы OpenCL v1.1 и примеры кода SDK доступны зарегистрированным разработчикам в области GPU вычислений. Скачайте OpenCL v1.1 уже сегодня – заходите на сайт или получите учетную запись.

Но с техникой OpenCL если и разбираться ... то уже в другой отдельной теме

[Olej]

Первый руками собственными сделанный многопроцессорный проект... для тестирования "как это происходит".
Элементарное сложение 2-х длинных векторов, действие элементарное, но это уже реальные потребности из области цифровой обработки сигналов.
Программный код я здесь рисовать не стану - его достаточно много, для интересующихся прилагается архив.

Результаты интересно наблюдать, в зависимости от длины векторов (ключ -n)

Код: Выделить всё

[olej@nvidia first]$ ./second -n 10000
GPU execution
A: 0.999758, B: 0.999881
C=A+B: 1.99977
time was 0.46314 million proc. cycles

[olej@nvidia first]$ ./second -n 10000 -c
CPU execution
A: 0.999758, B: 0.999881
C=A+B: 1.99977
time was 0.137995 million proc. cycles

[olej@nvidia first]$ ./second -n 100000
GPU execution
A: 0.999968, B: 0.999987
C=A+B: 1.9999
time was 0.50721 million proc. cycles

[olej@nvidia first]$ ./second -n 100000 -c
CPU execution
A: 0.999968, B: 0.999987
C=A+B: 1.9999
time was 1.346149 million proc. cycles

- при 10000 отсчётов CPU (ключ -c) ещё обгоняет GPU в 3 раза, но при 100000 векторах он уже отстаёт в 3 раза...
И это совсем не заурядный CPU:

Код: Выделить всё

[olej@nvidia first]$ ./testclc
wait ...
CPU speed: 3069412502 Hz

- это тест runtime измерения тактовой частоты процессора (из того же архива) - > 3GHz.
Т.е. можно считать, что на этих операциях 3GHz компьютер заработал как 10GHz

И это на элементарном действии, когда нужно просто поэлементно сложить массивы... на более сложных действиях разрыв, естественно, будет только нарастать.

[Olej]

Первый руками собственными сделанный многопроцессорный проект... для тестирования "как это происходит".

Сложность и неприятность программирования для CUDA, как оказалось, в том, что все существующие книжки, которые я перечислил в начале - полное барахло, об этом говорят и отзывы на них по интернету (цитировать не хочется чего я начитался) ... характерно, что ни в одной из них не сказано толком даже как толком скомпилировать такую программу компилятором nvcc, а приводимые в этих книгах примеры - не работоспособные!
А других нет...

Кроме, собственно, фирменного NVIDIA руководства CUDA_C_Programming_Guide.pdf, из комплекта SDK, которое тоже компактное, на 185 стр. с предисловиями и оглавлениями, и не покрывает всех вопросов... А дальше, получается: методом проб и ошибок.

[Olej]

Кроме, собственно, фирменного NVIDIA руководства CUDA_C_Programming_Guide.pdf, из комплекта SDK, которое тоже компактное, на 185 стр. с предисловиями и оглавлениями, и не покрывает всех вопросов... А дальше, получается: методом проб и ошибок.

Поиском нашёл ещё интересный (!) ресурс русскоязычный по CUDA (а если точнее, то не только CUDA, но и AMD ATI и OpenCL, т.е. использование сверхпроизводительных GPU для вычислений вообще):
http://www.gpgpu.ru/node/71
GPGPU.ru
Использование видеокарт для вычислений
Там же форум:
http://www.gpgpu.ru/forum

Поверхностное разбирательство с состоянием ... проб и ошибок по GPU в русскоязычном пространстве (обсуждения, блоги...) показало ещё одну любопытную вещь:
- значительная часть народа устанавливает CUDA SDK для экспериментов под Windows & Visual C...
- но как раз потенциальные возможности, сферы применений GPU вычислений сверхвысокой производительности - радикально отличаются под Windows & Linux ...
- под Linux это могут быть: системы реального времени (ну, почти реального времени ) радиолокации, гидролокации, какие-то задачи баллистики ... т.е. укрупнённо: из области "промышленной автоматики";
- под Windows это, скорее всего, различные системы online обработки видеопотоков, линейный монтаж etc.
- ну, и как мне кажется, как раз отработка GPU вычислений в Linux - интереснее, актуальнее...

[Olej]

Чтобы двигаться дальше с написанием проектов CUDA нужно решить возможность собирать такое приложение:
- из нескольких исходных файлов кода...
- которые компилируются отдельно (иногда это бывает разными компиляторами и с разных языков программирования)...
- а потом это всё (объектные файлы) собирается в единое приложение линковщиком ld, но обычно для этого указывают тот же gcc, который сам уже вызовет ld...

Так вот на этом деле CUDA мне крови попила! ... 2 дня убил на выкручивание ей рук...

Чтоб не объяснять множество-множество сообщений об ошибках (если кто пойдёт следом повторять), основные пункты:
- компилятор nvcc (NVIDIA компилятор для CUDA) своим фронт-эндом использует (в Linux) компилятор gcc, но для языковых стандартов С++ ...
- а компиляция с C & C++ по-разному создают внешние имена для связывания в объектных модулях .... поэтому дополнительные модули нужно компилировать в стандартах C++ (а не C);
- и nvcc и gcc при дефаултной команде используют умолчания для поиска своих разделяемых библиотек ... при сборке из объектных модулей эти умолчания (некоторых библиотек) придётся указать явно...

И чтоб не объяснять далее на пальцах, я нарисую Makefile как результат 3-го дня войны с nvcc:

Код: Выделить всё

ALL      = solid tstcc multiple

all:    $(ALL)

solid:  solid.cu clock.h clock.c
        nvcc solid.cu -o solid

multiple: multiple.cu clock.o
        nvcc --compile multiple.cu -o multiple.o
        g++ multiple.o clock.o -L/usr/local/cuda/lib/ -lcudart -o multiple

tstcc:  tstcc.cc clock.o
        g++ -c tstcc.cc -o tstcc.o
        g++ tstcc.o clock.o -o tstcc

clock.o: clock.h clock.c
        g++ -c clock.c -o clock.o

clean:
        rm -rf $(ALL) *.o

Здесь solid - это цельное приложение CUDA собираемое из одного файла, а multiple - в точности то же приложение, но разорванное на 2 исходных файла: multiple.cu - CUDA часть + clock.c - это чисто C код (который компилируется как C++ ).

[Olej]

Здесь solid - это цельное приложение CUDA собираемое из одного файла, а multiple - в точности то же приложение, но разорванное на 2 исходных файла: multiple.cu - CUDA часть + clock.c - это чисто C код (который компилируется как C++ ).

Код:
- solid.cu

Код: Выделить всё

#include <iostream>
using namespace std;
#include "clock.h"

#include "clock.c"  // реализация clock_cycles()

// Kernel definition
__global__ void VecAdd( void ) {
    return;
}

int main() {
    int N = 100;
    uint64_t t1, t2;
    t1 = clock_cycles();
    // Kernel invocation with N threads
    VecAdd<<<1, N>>>();
    t2 = clock_cycles();
    cout << "execution time was " << ( t2 - t1 ) << " cycles" << endl;
}

- и он текстуально включает (#include "clock.c") clock.c:

Код: Выделить всё

#include "clock.h"

uint64_t clock_cycles( void ) {
    uint64_t x;
    asm volatile ( "rdtsc" : "=A" (x) );
    return x;
}

- и альтернатива solid.cu - файл multiple.cu :

Код: Выделить всё

#include <iostream>
using namespace std;
#include "clock.h"

// Kernel definition
__global__ void VecAdd( void ) {
    return;
}

int main() {
    int N = 100;
    uint64_t t1, t2;
    t1 = clock_cycles();
    // Kernel invocation with N threads
    VecAdd<<<1, N>>>();
    t2 = clock_cycles();
    cout << "execution time was " << ( t2 - t1 ) << " cycles" << endl;
}

Вы скажете: "а чем он отличается?". А тем и отличается , то не включает ту единственную строчку #include "clock.c", а файл clock.c тупо компилируется отдельно (g++ не gcc !!!), и компонуются потом они вместе.

[Olej]

Код:

Этот простенький пример (упрощённый "до-нильзя") порождает много вопросов (больше чем даёт ответов) - а это основная техника сборки сложных реальных проектов (в отличие от тех примеров, что NVIDIA приводит в примеры CUDA):

- компиляция C-кода делается исключительно g++, хотя бы просто потому, что C & C++ образуют по-разному внешние имена для линковки:

Код: Выделить всё

[olej@nvidia mobj]$ nm clock.o
00000000 T _Z12clock_cyclesv

вот то _Z12... и есть типизация параметров и возврата C++.
Можно ли прикомпоновать объектный модуль скомпилированный C (gcc)? Не знаю, почти уверен, что да, ... обложив определение функций указанием "что это C", но это нужно проверять.

- стоит поменять g++ на gcc на линковке и всё рассыпается ... из-за дефаултных подключаемых библиотек (а как их дефаулт - посмотреть? я так с ходу не помню...):

Код: Выделить всё

[olej@nvidia mobj]$ ldd multiple
    linux-gate.so.1 =>  (0x0076b000)
    libcudart.so.4 => /usr/local/cuda/lib/libcudart.so.4 (0x00ddc000)
    libstdc++.so.6 => /usr/lib/libstdc++.so.6 (0x075d5000)
    libm.so.6 => /lib/libm.so.6 (0x00b51000)
    libgcc_s.so.1 => /lib/libgcc_s.so.1 (0x00b88000)
    libc.so.6 => /lib/libc.so.6 (0x009a1000)
    libdl.so.2 => /lib/libdl.so.2 (0x00b4a000)
    libpthread.so.0 => /lib/libpthread.so.0 (0x00b2d000)
    librt.so.1 => /lib/librt.so.1 (0x00b7d000)
    /lib/ld-linux.so.2 (0x0097c000)

- ну и в подтверждение того, что это одно и то же приложение (почти - размеры то чуть-чуть отличаются )

Код: Выделить всё

[olej@nvidia mobj]$ ls -l *
...
-rwxrwxr-x. 1 olej olej 11550 Ноя  5 21:09 multiple
...
-rwxrwxr-x. 1 olej olej 11525 Ноя  5 21:09 solid
...

вот выполнения полученного:

Код: Выделить всё

 
[olej@nvidia mobj]$ ./multiple
execution time was 184467452 cycles
[olej@nvidia mobj]$ ./solid
execution time was 132999604 cycles
[olej@nvidia mobj]$ ./solid
execution time was 160662372 cycles
[olej@nvidia mobj]$ ./multiple
execution time was 117963975 cycles

- пустое обращение к библиотекам CUDA и запуск пустых потоков выполнения потребовало около 150 миллионов тактов CPU.

[Olej]

Можно ли прикомпоновать объектный модуль скомпилированный C (gcc)? Не знаю, почти уверен, что да, ... обложив определение функций указанием "что это C", но это нужно проверять.

Это легко проверить (читаем Бьерн Страуструп "... Специальное издание", 2002г. - стр.250):

- в предыдущем проекте добавляем 3 строчки в файл clock.h:

Код: Выделить всё

#include <stdio.h>
#include <stdlib.h>
#include <stdint.h>
#include <unistd.h>
#include <time.h>

#ifdef CONLY
extern "C"
#endif
uint64_t clock_cycles( void );

это те 3 строки перед uint64_t ...
Всё! этого достаточно для проверки. Но сборку делаем теперь по-другому (чуть-чуть):

Код: Выделить всё

[olej@nvidia mobj]$ g++ -D CONLY -c clock.c -o clock_c.o
[olej@nvidia mobj]$ nm clock_c.o
00000000 T clock_cycles
[olej@nvidia mobj]$ nvcc --compile -Xcompiler '-D CONLY'  multiple.cu -o multiple_с.o
[olej@nvidia mobj]$ nm multiple_с.o | grep clock
         U clock_cycles

- глядите, как поменялось имя связывания: теперь это классическое C имя, никаких соглашений C++ ... такое имя может прийти из любого отдельно скомпилированного gcc объектного модуля.
Имя поменялось в обоих модулях синхронно ... самое рискованное здесь было "научить" nvcc

Код: Выделить всё

[olej@nvidia mobj]$ g++ multiple_с.o clock_c.o -L/usr/local/cuda/lib/ -lcudart -o multiple
[olej@nvidia mobj]$ ./multiple
execution time was 151878154 cycles

Всё, и этот вопрос положительно разрешён.

[Olej]

http://www.gpgpu.ru/reviews

October 20th, 2011
Последние дни оказались богаты на всякие штуки для OpenCL:

OpenCL Compiler tools
Небольшая библиотека, позволяющая удобно компилировать OpenCL kernels заранее, сохраняя результат (можно - шифрованый). Цель - поэкономить на времени компиляции и на том, что kernel находится где-то внутри программы в исходных текстах.
Несколько смазывает впечатление то, что библиотека распространяется под GPL3, т.е. цель "спрятать kernel" под такой лицензией, очевидно, не может быть достигнута т.к. использующий библиотеку варез обязан тоже быть под GPL3+

Насколько всё-таки отличаются ... философии в зависимости от используемых tools :
- это цитата с очень неплохого сайта, где занимаются отработкой CUDA (по-моему, это МГУ или около)...
- но занимаются на Windows...
- и их беспокоит, что у них не будет средств закрыть своё ПО ... которого ещё далеко нет в природе, и неизвестно, будет ли и что...
- в то время, когда в Linux сильно переживают: не будет ли лицензия GPL3 слишком ограничивать открытость, доступность и фришность...