Сложение векторов и сравнение с количеством тредов в памяти
Данный пример довольно прост и демонстрирует реализацию сложения двух векторов, у которых сумма общих нитей не превышает определенное число (в данном случае 10).
#include <iostream> #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; }
Понятное дело, если мы будем использовать N Блоков и 1 тред (add<<< N, 1 >>>( dev_a, dev_b, dev_c ); ), то каждый результат с будет равным нулю, кроме первого, поэтому этот код закомментирован, но присутствует в программе в качестве ознакомления с работой блоков и тредов в ядре. В нашем случае результат выполнения программы следующий:
И так до 63 значений, так как подсчет идет с нуля.
Генерация псевдослучайных чисел с использованием CuRand
Генерация псевдослучайных чисел происходит с помощью инструмента curand, который входит в расширение CUDA C.
Библиотека cuRAND предоставляет средства, которые сосредоточены на простой и эффективной генерации псевдослучайных и квазислучайных чисел. cuRAND состоит из двух частей: библиотеки на стороне хоста (CPU) и файла заголовка устройства (GPU). Библиотека на стороне хоста рассматривается как любая другая библиотека: пользователь включает заголовочный файл, #include<curand.h>, чтобы получить объявления функций. Случайные числа могут быть сгенерированы на девайсе или на хосте. Полученные случайные числа сохраняются в глобальной памяти на девайсе. Также пользователь может копировать полученные случайные числа обратно на хост для дальнейшей обработки. Для генерации чисел на хосте вся работа выполняется на хосте, включая их хранение.
Вторая составляющая cuRAND — это заголовочный файл устройства, /include/curand_kernel.h. Этот файл определяет функции генератора случайных чисел. Пользовательский код может включать этот заголовочный файл, а написанные пользователем ядра могут затем вызывать функции устройства, определенные в файле заголовка. Это позволяет создавать случайные числа и немедленно использовать их.
Ниже будет код, демонстрирующий работу данного инструмента:
#include <curand.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, а также здесь происходит работа с памятью – выделение и освобождение. Результатом выполнения программы будет генерация случайного числа:
