Подскажите, где можно почитать по этой теме? Какие основные идеи, в чем разница по сравнению с программированием CPU? Каковы ограничения кода, как устроена память и взаимодействие с хостом, где основные бутылочные горлышки?
Здравствуйте, Tekk, Вы писали:
T>Подскажите, где можно почитать по этой теме? Какие основные идеи, в чем разница по сравнению с программированием CPU? Каковы ограничения кода, как устроена память и взаимодействие с хостом, где основные бутылочные горлышки?
Я несколько месяцев назад поставил себе CUDA, покурил примеры, нашёл русскоязычные статьи, сам написал простенькие программки. Узкие места неплохо показывает профайлер.
После реализовал свой детектор движения, тот ускорился в 4 раза. Причём больше всего времени занимает копирование кадров в видеопамять. На этом довольный спокоился.
Здравствуйте, Nuzhny, Вы писали:
N>Я несколько месяцев назад поставил себе CUDA, покурил примеры, нашёл русскоязычные статьи, сам написал простенькие программки. Узкие места неплохо показывает профайлер.
Насколько я понимаю, сейчас перспективнее всего OpenCL. Разве нет?
T>Насколько я понимаю, сейчас перспективнее всего OpenCL. Разве нет?
Да их как зайцев, ещё забыли ati stream и dx11 computing shaders. И larrabee — только он умер до рождения.
Здравствуйте, Tekk, Вы писали:
T>OpenCL — кросс-платформа. В отличие от куды, например.
Дык не спорю. Скорее всего она и останется. Ну и dx11 разумеется
Здравствуйте, Tekk, Вы писали:
N>>Я несколько месяцев назад поставил себе CUDA, покурил примеры, нашёл русскоязычные статьи, сам написал простенькие программки. Узкие места неплохо показывает профайлер.
T>Насколько я понимаю, сейчас перспективнее всего OpenCL. Разве нет?
Может быть. Однако у нас специфическое ПО, которое чаще всего поставляется вместе с железом, поэтому пока можно позволить себе ограничиться видеокартами nVidia. Надо посмотреть на востребованность, если что перепишу.
Главный критерий при выборе, который передо мной тогда стоял, — быстрое освоение технологии. Буквально неделя у меня ушла на чтение документации, статей и опыты, а после неделя ушла на портирование самого модуля, отладку и профилирование. В результате четырёхкратное ускрение на разрешении 768х576 и двукратное на меньших.
Здравствуйте, Tekk, Вы писали:
T> Какие основные идеи, в чем разница по сравнению с программированием CPU? Каковы ограничения кода, как устроена память и взаимодействие с хостом, где основные бутылочные горлышки?
ИМХО,
Основное отличие от CPU — то что GPU подходит в основном только для векторных алгоритмов. А там где много ветвящихся if-else, GPU не очень подходит.
В идеале, лучшая производительность когда одна и та-же функция без if-else блоков запускается на множестве ядер, и разные экземпляры функций отличаются индексами, по которым можно обращаться к разным участкам памяти .
Даже физическая реализация такая. У GeForce карточек ядра сгруппированы по 8 штук в мультипроцессор. Если написано 128 ядер, значит только 16 независимых мультипроцессоров. Да и то их затруднительно загрузить разными задачами. А внутри мультипроцессора — 8 ядер могут только одинаковые для всех команды исполнять, имеют только разные индексы(чтоб из них разные адреса вычислять). Ветвления if,else фактически не возможны. Хотя формально есть if-else, но если узнаешь как они выполняются, то использовать их не захочется.
Быстрая память для доступа из ядер маленькая — пара десятков килобайт. А при доступе векторных функций к медленой памяти надо тщательно выравнивать, шаманские трюки применять, иначе на порядок медленнее будет работать.
Все достаточно низкоуровневое. И плохо что много неявных особенностей в плане перфоманса , которые надо проверять методом тыка.
По Compute Shaders в DX11 информации мало, в SDK только 2 примера, копирование массива и сортировка. В Cuda SDK побольше можно найти.
Ради ускорения в 2 раза, возиться не интересно. Разве что, стоит повозиться если удатся сделать ускорение пропорционально числу ядер. Но это получится только для некоторых алгоритмов.
Здравствуйте, Silver_s, Вы писали:
S_>Ветвления if,else фактически не возможны. Хотя формально есть if-else, но если узнаешь как они выполняются, то использовать их не захочется.
А что там с ними, можно подробнее?
И еще интересно, что там с временем доступа. Локальная и глобальная память карточки, перекачка данных из оперативки — как они примерно соотносятся по скорости?
Здравствуйте, Tekk, Вы писали: T>Здравствуйте, Silver_s, Вы писали: S_>>Ветвления if,else фактически не возможны. Хотя формально есть if-else, но если узнаешь как они выполняются, то использовать их не захочется.
T>А что там с ними, можно подробнее? T>И еще интересно, что там с временем доступа. Локальная и глобальная память карточки, перекачка данных из оперативки — как они примерно соотносятся по скорости?
Все описать достоверно и исчерпывающе не смогу.
Но в двух словах. Если в множестве связанных потоков/ядер происходит ветвление, тогда все останавливаются и ждут пока ответвившиеся не сделают свое дело, и потом все вместе двигаются дальше. Каждая очередная команда на всех общая. Если много ветвлений то в этом связанном множестве просто утрачивается паралельность. И этот мультипроцессор из 8 ядер, будет по производительности ближе к одному ядру.
Скину некторые обрезки которые для себя надергал из Cuda programmer's guide, может там есть ответы на некоторые вопросы:
-A multiprocessor consists of eight Scalar Processor (SP) cores, two special function units for transcendentals,
a multithreaded instruction unit, and on-chip shared memory.
— The multiprocessor SIMT unit creates, manages, schedules, and executes threads in groups of 32 parallel threads called warps.
— If threads of a warp diverge via a datadependent conditional branch, the warp serially executes each branch path taken,
disabling threads that are not on that path, and when all paths complete, the threads converge back to the same execution path.
Branch divergence occurs only within a warp; different warps execute independently regardless of whether they are executing common or disjointed code paths.
- Each thread has a private local memory.
— Each thread block has a shared memory visible to all threads of the block and with the same lifetime as the block.
Finally, all threads have access to the same global memory.
-There are also two additional read-only memory spaces accessible by all threads: the constant and texture memory spaces.
-The global, constant, and texture memory spaces are persistent across kernel launches by the same application
-Shared memory is expected to be much faster than global memory as mentioned in Section 2.2 and detailed in Section 5.1.2.5.
-----------
Throughput of memory operations is 8 operations per clock cycle.
When accessing local or global memory, there are, in addition, 400 to 600 clock cycles of memory latency.
As an example, the throughput for the assignment operator in the following sample code:
__shared__ float shared[32];
__device__ float device[32];
shared[threadIdx.x] = device[threadIdx.x];
is 8 operations per clock cycle for the read from global memory, 8 operations per clock cycle for the write to shared memory,
but above all, there is a latency of 400 to 600 clock cycles to read data from global memory.
------------------
Global Memory
The global memory space is not cached, so it is all the more important to follow the right access pattern to get maximum memory bandwidth,
especially given how costly accesses to device memory are.
First, the device is capable of reading 4-byte, 8-byte, or 16-byte words from global memory into registers in a single instruction.
To have assignments such as:
__device__ type device[32];
type data = device[tid];
compile to a single load instruction, type must be such that sizeof(type) is equal to 4, 8, or 16 and variables of type type must be aligned to sizeof(type)
bytes (that is, have their address be a multiple of sizeof(type)).
-------------------
Shared Memory
Because it is on-chip, the shared memory space is much faster than the local and global memory spaces.
In fact, for all threads of a warp, accessing the shared memory is as fast as accessing a register as long as there are no bank conflicts between the threads, as detailed below.
--------------------
Registers
Generally, accessing a register is zero extra clock cycles per instruction,
but delays may occur due to register read-after-write dependencies and register memory bank conflicts.
------------------
A typical programming pattern is to stage data coming from device memory into shared memory;
in other words, to have each thread of a block:
— Load data from device memory to shared memory,
— Synchronize with all the other threads of the block so that each thread can safely read shared memory locations that were written by different threads,
— Process the data in shared memory,
— Synchronize again if necessary to make sure that shared memory has been updated with the results,
— Write the results back to device memory.
First, the device is capable of reading 4-byte, 8-byte, or 16-byte words from global memory into registers in a single instruction.
To have assignments such as:
__device__ type device[32];
type data = device[tid];
Specifications for Compute Capability 1.0
— The maximum number of threads per block is 512;
— The maximum sizes of the x-, y-, and z-dimension of a thread block are 512, 512, and 64, respectively;
— The maximum size of each dimension of a grid of thread blocks is 65535;
— The warp size is 32 threads;
— The number of registers per multiprocessor is 8192;
— The amount of shared memory available per multiprocessor is 16 KB organized into 16 banks (see Section 5.1.2.5);
— The total amount of constant memory is 64 KB;
— The total amount of local memory per thread is 16 KB;
— The cache working set for constant memory is 8 KB per multiprocessor;
— The cache working set for texture memory varies between 6 and 8 KB per multiprocessor;
— The maximum number of active blocks per multiprocessor is 8;
— The maximum number of active warps per multiprocessor is 24;
— The maximum number of active threads per multiprocessor is 768;
— For a one-dimensional texture reference bound to a CUDA array, the maximum width is 213;
— For a one-dimensional texture reference bound to linear memory, the maximum width is 227;
— For a two-dimensional texture reference bound to linear memory or a CUDA array, the maximum width is 216 and the maximum height is 215;
— For a three-dimensional texture reference bound to a CUDA array, the maximum width is 211, the maximum height is 211,
and the maximum depth is211;
— The limit on kernel size is 2 millions of microcode instructions;
Здравствуйте, Tekk, Вы писали:
T>Подскажите, где можно почитать по этой теме? Какие основные идеи, в чем разница по сравнению с программированием CPU? Каковы ограничения кода, как устроена память и взаимодействие с хостом, где основные бутылочные горлышки?
У меня карта Radeon hd3870 и основное бутылочное горлышко, которое сводит на нет все программирование — медленное чтение данных ИЗ карты.
Жалкие 100Мб/с. Хотя запись — вполне терпимые 500Мб/с.
Проверял в windows XP 32bit, windows 7 32 bit и ubuntu linux 64 bit — везде одинаково.
Здравствуйте, Sinix, Вы писали:
S>Здравствуйте, Tekk, Вы писали:
T>>OpenCL — кросс-платформа. В отличие от куды, например. S>Дык не спорю. Скорее всего она и останется. Ну и dx11 разумеется
S>Бенчмарков всей компании не попадалось?
нет, но один и тот же пример из NVIDIA GPU computing SDK на OpenCL работает на глаз заметно медленнее, чем на CUDA C. Так что оно конечно перспективное, но пока реализации сыроваты.
Здравствуйте, Sergey Chadov!
SC>нет, но один и тот же пример из NVIDIA GPU computing SDK на OpenCL работает на глаз заметно медленнее, чем на CUDA C. Так что оно конечно перспективное, но пока реализации сыроваты.
The latest conformant release driver/OpenCL run-time brings performance parity with CUDA, a great improvement on the preview release. The public release (due soon) might even be faster than CUDA and can only get better. There is no reason not to port CUDA code to OpenCL now!
N>В результате четырёхкратное ускрение на разрешении 768х576 и двукратное на меньших.
в рамках своего текущего проекта, я реализовал код работающий в ~1400 раз (пик) быстрее аналога на цпу.
конечно задача попалась такая, что позволяет писать код такой производительности.
Здравствуйте, Muxa, Вы писали:
M>в рамках своего текущего проекта, я реализовал код работающий в ~1400 раз (пик) быстрее аналога на цпу. M>конечно задача попалась такая, что позволяет писать код такой производительности.
Это просто сказочные цифры. Если не секрет, то что за задача (или хотя бы область)?
Замеры проводились с учётом копирования исходных данных в память карты и результата в системную?
N>Это просто сказочные цифры. Если не секрет, то что за задача (или хотя бы область)?
см. подпись.
N>Замеры проводились с учётом копирования исходных данных в память карты и результата в системную?
учитывалось все.
