Занятие 8. Программирование на CUDA C. Часть 4.

Сложение векторов и сравнение с количеством тредов в памяти

Данный пример довольно прост и демонстрирует реализацию сложения двух векторов, у которых сумма общих нитей не превышает определенное число (в данном случае 10).

#include <cuda.h>
 
using namespace std;
 
__global__ void add( float *a, float *b, float *c ) {
if(a[ threadIdx.x ] + b[ threadIdx.x ]<10)
c[ threadIdx.x ] = a[ threadIdx.x ] + b[ threadIdx.x ];
else
c[ threadIdx.x ] = 10;
}
 
#define N 64 //Количество суммирований и сравнений
 
int main( void ) {
float *a, *b, *c; // host копии a, b, c
float *dev_a, *dev_b, *dev_c; // device копии a, b, c
int size = N * sizeof( float );
 
//выделяем память для device копий a, b, c
 
cudaMalloc( (void**)&dev_a, size );
cudaMalloc( (void**)&dev_b, size );
cudaMalloc( (void**)&dev_c, size );
 
a = (float*)malloc( size );
b = (float*)malloc( size );
c = (float*)malloc( size );
 
for (int i = 0; i < N; ++i)
a[i] = (float)rand()/(float)RAND_MAX;
 
for (int i = 0; i < N; ++i)
b[i] = (float)rand()/(float)RAND_MAX;
 
// копируем ввод на device
 
cudaMemcpy( dev_a, a, size, cudaMemcpyHostToDevice );
cudaMemcpy( dev_b, b, size, cudaMemcpyHostToDevice );
 
// launch add() kernel with N parallel blocks
 
add<<< 1, N >>>( dev_a, dev_b, dev_c ); //Использование N тредов и 1 блока
 
// add<<< N, 1 >>>( dev_a, dev_b, dev_c ); //Использование N блоков И 1 треда
 
// копируем результат работы device обратно на host - копию c
 
cudaMemcpy( c, dev_c, size, cudaMemcpyDeviceToHost );
 
for(int i=0; i<N; i++)
 
cout<<"c["<<i<<"]="<<c[i]<<endl;
 
free( a ); free( b ); free( c );
cudaFree( dev_a );
cudaFree( dev_b );
cudaFree( dev_c );
 
return 0;
 
}

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

И так до 63 значений, так как подсчет идет с нуля.

Генерация псевдослучайных чисел с использованием CuRand

Генерация псевдослучайных чисел происходит с помощью инструмента curand, который входит в расширение CUDA C.

, чтобы получить объявления функций. Случайные числа могут быть сгенерированы на девайсе или на хосте. Полученные случайные числа сохраняются в глобальной памяти на девайсе. Также пользователь может копировать полученные случайные числа обратно на хост для дальнейшей обработки. Для генерации чисел на хосте вся работа выполняется на хосте, включая их хранение.

/include/curand_kernel.h. Этот файл определяет функции генератора случайных чисел. Пользовательский код может включать этот заголовочный файл, а написанные пользователем ядра могут затем вызывать функции устройства, определенные в файле заголовка. Это позволяет создавать случайные числа и немедленно использовать их.

Ниже будет код, демонстрирующий работу данного инструмента:

#include <curand_kernel.h>
 
#define MAX 100
 
/* эта функция ядра GPU вычисляет случайное число и сохраняет его в памяти*/
 
__global__ void random(unsigned int seed, int* result) {
 
/* Библиотека случайных чисел CUDA использует curandState_t для отслеживания начального значения мы будем хранить случайное состояние для каждого потока*/
 
curandState_t state;
 
/* инициализация состояния*/
 
curand_init(seed, /* seed контролирует последовательность значений, которые генерируются*/
 
0, /* порядковый номер важен только с несколькими ядрами*/
 
0,
 
&state);
 
/* curand работает как rand - за исключением того, что он принимает состояние как параметр*/
 
*result = curand(&state) % MAX;
}
int main( ) {
 
/* выделить память int на GPU*/
 
int* gpu_x;
cudaMalloc((void**) &gpu_x, sizeof(int));
 
/* вызывать GPU для инициализации всех случайных состояний*/
 
random<<<1, 1>>>(time(NULL), gpu_x);
 
/* скопировать случайное число на CPU*/
 
int x;
cudaMemcpy(&x, gpu_x, sizeof(int), cudaMemcpyDeviceToHost);
printf("Random number = %d.\n", x);
 
/* освобождение памяти */
 
cudaFree(gpu_x);
 
return 0;

Как было сказано, первоначально объявляются заголовочные файлы, после чего прописывается стандартная С++ директива define с макросом MAX 100, обозначающая, что случайное число будет генерироваться не больше 100. Затем происходит вызов функции на GPU, генерирующей случайное число. Значение будет записываться в переменную result. В главной функции (main) происходит инициализация всех значений через вызов GPU, а также здесь происходит работа с памятью – выделение и освобождение. Результатом выполнения программы будет генерация случайного числа: