Основной плюс использования CUDA это параллелизм. Представьте можно брутить хэш в несколько поток с огромной скоростью, которой обычным процессорам пока и не снилась. Например мой 9600(причем один из самых дешевых) содержит 6 мультипроцессоров, в каждом 8-10 ядер и сотней ALU, несколькими тысячами регистров. Ядра могут исполнять один и тот же код одновременно. Видеочипы могут поддерживать до 1024 потоков на каждый мультипроцессор, при том переключаться между потоками можно несколько раз за один такт.

Немного о архитектуре NVIDIA. GPU(графический процессор) состоит из нескольких функциональных блоков Texture Processing Clusters (кластерных текстурных блоков), они состоят из нескольких Streaming Multiprocessor (потоковых мультипроцессоров) представляющих собой 8 вычислительных ядер Streaming Processors (потоковых процессоров) и блоков адресации и фильтрации текстур. Все инструкции выполняются по принципу SIMD(Single Instruction, Multiple Data – одна инструкция для множества данных, процессор состоит из нескольких модулей обработки данных(процессорных элементов) и командного процессора(управляющего модуля, контроллера). Если в команде поступившей на процессор встречаются данные, процессор рассылает ее на все элементы и она там выполняется , каждый процессорный элемент имеет свою собственную память), когда одна инструкция применяется ко всем нитям в warp(группа из 32 нитей – минимальный объем данных, обрабатываемый одним мультипроцессором).

Каждый мультипроцессор имеет определенные ресурсы: разделяемую память (используются потоками одного блока, нити из разных блоков обмениваться данными не могут), кэш на константы и текстуры.

Программная модель CUDA

Все нити группируются в иерархию – grid/block/thread.

Самый верхний уровень это grid объединяет все нити выполняющие ядро, grid одно-мерный или двухмерный массив блоков. Block(блок) одномерный или двухмерный или трехмерный массив нитей.

Итак, kernel(ядро) – это функция, которая выполняется на GPU. Что бы объявить та-кую функцию надо воспользоваться спецификатором __global__. Что бы вызвать эту функцию надо пользоваться специальным синтаксисом, тройными угловыми скобками, синтаксис такой:

kernelName<<<Dg,Db,Ns,S>>>(args)

Отступление: здесь стоит упомянуть о типах данных(1,2,3,4-мерные вектора) используемых в API CUDA. Они описываются в хедере vector_types.h как структуры. Для их создания используется вызов функции вида make_<type_name>(args), например, для создания вектора uint3 пишем: uint3 u3=make_uint3(1,2,3), (посмотрите хедер vector_function.h). 

Dg – переменная типа dim3, задающий размерность и размер grid в блоках. Например если dim3 grid=dim3(16,16), то грид двухмерный, состоящий из 16Х16X1=256 блоков. Здесь один момент, запустите программку Device Query Driver API из CUDA SDK, она покажет некоторые характеристика вашей карточки. Так вот там есть такая характеристи-ка cudaDeviceProp.maxGridSize[3] показывающая максимальный размер каждой компоненты grid, у меня например это 65535Х65535Х1, т.е. grid двумерный.

Db – переменная типа dim3, задающий размерность и размер блока в нитях. Через cudaDeviceProp.maxThreadsDim[3] можно узнать максимальный размер каждой компонен-ты блока, у меня это 512Х512Х64(максимальная размерность блока 3), есть еще одна характеристика карточки cudaDeviceProp.maxThreadsPerBlock описывающая максимальное число нитей в блоке, у меня 512, следовательно, можно использовать блок размером и размерностью, например 16Х16Х2, или 256Х2, или 512 (что количество нитей было максимум 512).

Ns – переменная типа size_t задает дополнительный объем shared памяти, которая должна быть динамически выделена (к уже статически выделенной). Характеристика cudaDeviceProp.sharedMemPerBlock показывает максимум shared памяти на блок в байтах, у меня 16384. По умолчанию этот параметр равен 0.

S – задает поток в котором должен произойти вызов, по умолчанию равен 0.

__global__ void VecAdd(float* A, float* B, float* C)

{

...

}

int main()

{

...

VecAdd<<<1, N>>>(A, B, C);

}

VecAdd<<<1, N>>>(A, B, C) значит мы вызвали kernel VecAdd с Dg равным 1X1X1, Db Nx1x1, Ns=0, S=0, говоря проще одномерный грид содержащий одномерный блок, который в свою очередь содержит N нитей. 

Каждая нить выполняющая ядро имеет свой уникальный ID, доступный в ядре через переменную threadIdx. Ниже пример сложения двух векторов размерностью N, и присвоения результата вектору C.

__global__ void VecAdd(float* A, float* B, float* C)

{

int i = threadIdx.x;

C[i] = A[i] + B[i];

}

int main()

{

...

VecAdd<<<1, N>>>(A, B, C);

}

Вектора А,В,С имеют размерность N и используем мы для получения вектора C N нитей, т. е. каждая нить рассчитывает одну компоненту вектора С. Доступ к компонентам осуществляется через ID нити. threadIdx – это трехкомпонентный вектор(а точнее эта пе-ременная имеет тип uint3), указывающий адрес текущей нити в блоке, т.е. если через threadIdx мы можем получить доступ к одной компоненте, то блок одномерный, если к двум компонентам, блок двухмерный, к трем, трехмерный. 

Как найти индекс элемента массива, передаваемого в kernel,исходя из алгоритма описанного выше, для каждой нити? Если блок одномерный то все просто индекс равен threadIdx.x, если блок двухмерный, размер его (Dx,Dy), индекс рассчитывается по форму-ле (x + y Dx), для трехмерно блока (x + y Dx + z Dx Dy). Давайте рассмотрим несколько случаев для вычисления суммы двух N мерных векторов. 

1. Вызываем kernel с параметрами <<<1, N>>> получаем грид 1Х1Х1 и блок NX1X1, следовательно, как видно из примера выше индекс для доступ к элемен-ту массива А равен threadIdx.x

2. Вызываем kernel с параметрами <<<1, b>>>, где b=(b1,b2) и b1Xb2=N, тогда ин-декс равен threadIdx.x+ threadIdx.y*b1

Добавлены следующие зарезервированные переменные:

- gridDim – размер грида (dim3)

- blockDim – размер блока (dim3)

- blockIdx – индекс текущего блока в grid(uint3)

Нити в пределах блока могут использовать для обмена shared memory, так же они могут синхронизироваться. Добавлена встроенная функция _syncthread(), осуществляющая синхронизацию всех нитей блока. Эта функция действует как «барьер», пока все нити блока не выполнят ее они не смогут выполнять свой код дальше. Удобна для бесконфликтной работе с shared памятью.

Блоки выполняются независимо друг от друга, в любом порядке, параллельно или последовательно. Каждая нить имеет свою локальную память, каждый блок имеет свою shared(общую, совместную) память разделяемую всеми потоками блока и наконец все нити имеют доступ к глобальной памяти.

Еще существует два вида памяти доступных всем потокам только для чтения, это константная и текстурная память.

Модель CUDA предполагает что нити выполняются на device(GPU), который рабо-тает как сопроцессор к host(CPU). Ядро выполняется на GPU, а остальная программа вы-полняется на CPU

верстаем

Ссылки:

http://itc.ua/node/32663

http://www.ixbt.com/video3/cuda-1.shtml

http://ru.wikipedia.org/wiki/SIMD

http://developer.download.nvidia.com/compute/cuda/2_3/toolkit/docs/NVIDIA_CUDA_Programming_Guide_2.3.pdf