Скалярное произведение с использованием shared-памяти
Треды на первый взгляд кажутся избыточными. Параллельные треды, в отличие параллельных блоков, имеют механизмы для, так называемых, коммуникации и синхронизации. Давайте рассмотрим этот вопрос.
Возьмем в качестве примера на этот раз не сложение, а скалярное произведение.
Треды в одном блоке могут иметь общую память shared memory. Эта память декларируется в коде ключевым словом __shared__. Эта память не видна в другом блоке.
Также треды можно «синхронизировать». То есть каждый тред в некой «точке синхронизации» будет ждать пока выполнение других тредов тоже не достигнет этой точки. Синхронизация нитей истинна только для тредов одного блока. Синхронизация вводиться функцией __syncthreads().
Ниже представлен код с синхронизацией и использованием shared memory:
//__global__ void dot( int *a, int *b, int *c ) { // Каждый тред вычисляет одно слагаемое скалярного произведения // int temp = a[threadIdx.x] * b[threadIdx.x]; //Не может вычислить сумму //Каждая копия ‘temp’ приватна для своего треда //} __global__ void dot( int *a, int *b, int *c ) { __shared__ int temp[N]; temp[threadIdx.x] = a[threadIdx.x] * b[threadIdx.x]; __syncthreads(); if( 0 == threadIdx.x ) { int sum = 0; for( int i = 0; i < N; i++ ) sum += temp[i]; *c = sum; } } #define N 512 int main( void ) { int *a, *b, *c; int *dev_a, *dev_b, *dev_c; int size = N * sizeof( int ); cudaMalloc( (void**)&dev_a, size ); cudaMalloc( (void**)&dev_b, size ); cudaMalloc( (void**)&dev_c, sizeof( int ) ); a = (int *)malloc( size ); b = (int *)malloc( size ); c = (int *)malloc( sizeof( int ) ); random_ints( a, N ); random_ints( b, N ); // копируем ввод на device cudaMemcpy( dev_a, a, size, cudaMemcpyHostToDevice ); cudaMemcpy( dev_b, b, size, cudaMemcpyHostToDevice ); //запускаем на выполнение dot() kernel с 1 блоком и N тредами dot<<< 1, N >>>( dev_a, dev_b, dev_c ); //копируем результат работы device на host копией c cudaMemcpy( c, dev_c, sizeof( int ) , cudaMemcpyDeviceToHost ); free( a ); free( b ); free( c ); cudaFree( dev_a ); cudaFree( dev_b ); cudaFree( dev_c ); return 0; }
threadIdx.x используется для доступа к отдельной нити. threadIdx — это встроенная переменная. Для индексации вводимых/выводимых элементов используем выражение (threadIdx.x + blockIdx.x * blockDim.x).blockDim.x — это встроенная переменная, обозначающая число тредов в блоке.
Сложение двух векторов
При создании любого проекта на CUDA в Visual Studio происходит генерация кода, который демонстрирует сложение двух векторов. Это очень полезный и показательный пример работы с CUDA. Он будет представлен в следующем разделе, а пока давайте создадим собственную программу с нуля.
Необходимые библиотеки:
#include <cstdlib> #include <iostream> #include <cuda_runtime.h>
<cuda_runtime.h> является библиотекой от расширения CUDA C. Она нужна для того, чтобы можно было получать встроенные модули (например, типы данных) и была возможность использования API CUDA во время выполнения компиляции с хоста. Если вы будете использовать компилятор CUDA nvcc, то прописывать данную библиотеку необязательно, так как этот заголовок подключается автоматически (помимо cuda_runtime.h, это также касается cuda_runtime_api.h и cuda.h).
Следующим шагом будет добавление спецификатора __global__, его функция с аргументами. Среди аргументов обозначим 2 вектора const float * a и const float * b, результат float * result и счетчик int n:
__global__ void vecAdd_kernel( const float * a, const float * b, float * result, int n)
Далее введем новую переменную целочисленного типа i со значением blockIdx.x * blockDim.x + threadIdx.x, обозначающую размерность блоков и грида. В данном случае одномерная модель блоков и грида. В таком случае номер треда будет определяться как номер блока, умноженный на количество тредов в блоке, плюс номер треда внутри блока.
Дальше идет сравнение переменной i со счетчиком и запись результата сложения.
if (i < n) result[i] = a[i] + b[i];
После чего в главной функции программы задаем нашему счетчику значение 100, также в главной функции выделяем память под каждую переменную, включая полученный результат.
int n = 100; float * a = new float[n], * a_gpu; cudaMalloc((void**)&a_gpu, n * sizeof(float)); float * b = new float[n], * b_gpu; cudaMalloc((void**)&b_gpu, n * sizeof(float)); float * result = new float[n], * result_gpu; cudaMalloc((void**)&result_gpu, n * sizeof(float));
Добавляем цикл со счетчиком for с условиями i< n; i++, а в теле цикла будет сравниваться a и b с i при этом i будет увеличиваться на единицу (i++).
Далее происходит копирование данных из памяти CPU в память GPU с помощью функции cudaMemcpy. В аргументах указываем переменную с GPU и CPU, выделяемый размер и команду cudaMemcpyHostToDevice:
cudaMemcpy(a_gpu, a, n * sizeof(float), cudaMemcpyHostToDevice); cudaMemcpy(b_gpu, b, n * sizeof(float), cudaMemcpyHostToDevice);
Далее обозначаем размер блоков const int block_size = 256 и считаем количество блоков, которое потом будет вызываться в ядре:
int num_blocks = (n + block_size - 1) / block_size
Вызываем ядро, записываем результат и снова через cudaMemcpy передаем его уже с GPU на CPU (cudaMemcpyDeviceToHost):
vecAdd_kernel <<< num_blocks, block_size >>> (a_gpu, b_gpu, result_gpu, n); cudaMemcpy(result, result_gpu, n * sizeof(float), cudaMemcpyDeviceToHost);
Снова создаем цикл, который в дальнейшем напечатает результаты сложения векторов, которых будет равно количеству x:
for (int x = 0; x < 10; x++) cout<<result[x]<<endl;
Уничтожаем (освобождаем) массивы a, b, result на CPU и на GPU с помощью команды cudaFree:
delete [] a; delete [] b; delete [] result; cudaFree(a_gpu); cudaFree(b_gpu); cudaFree(result_gpu);
Таким образом, была реализована программа, которая складывает 2 вектора и показывает на экране результат их сложения. Листинг всей программы приведен ниже:
#include <cstdlib> #include <iostream> #include <cuda_runtime.h> using namespace std; __global__ void vecAdd_kernel( const float * a, const float * b, float * result, int n) { int i = blockIdx.x * blockDim.x + threadIdx.x; if (i < n) result[i] = a[i] + b[i]; } int main() { int n = 100; float * a = new float[n], * a_gpu; cudaMalloc((void**)&a_gpu, n * sizeof(float)); float * b = new float[n], * b_gpu; cudaMalloc((void**)&b_gpu, n * sizeof(float)); float * result = new float[n], * result_gpu; cudaMalloc((void**)&result_gpu, n * sizeof(float)); for (int i = 0; i < n; i++) a[i] = b[i] = i; cudaMemcpy(a_gpu, a, n * sizeof(float), cudaMemcpyHostToDevice); cudaMemcpy(b_gpu, b, n * sizeof(float), cudaMemcpyHostToDevice); const int block_size = 256; int num_blocks = (n + block_size - 1) / block_size; vecAdd_kernel <<< num_blocks, block_size >>> (a_gpu, b_gpu, result_gpu, n); cudaMemcpy(result, result_gpu, n * sizeof(float), cudaMemcpyDeviceToHost); for (int x = 0; x < 10; x++) cout<<result[x]<<endl; delete [] a; delete [] b; delete [] result; cudaFree(a_gpu); cudaFree(b_gpu); cudaFree(result_gpu); return 0; }
А результат выполнения программы вы можете увидеть на рисунке:
