Половинкин А.Н.

Постановка задачи Алгоритм вычисления функции axpy на GPU Программная реализация

y := alpha*x + y x, y – векторы размерность – n alpha – скаляр

BLOCK_SIZE... THREAD BLOCK 1 THREAD BLOCK 2 THREAD BLOCK K n Каждый блок потоков занимается вычислением одного подвектора y sub,i вектора y Каждый поток внутри блока потоков занимается вычислением одного элемента подвектора y sub,i BLOCK_SIZE THREAD BLOCK i... y sub,i

axpy.h – содержит определения (через define) размера блока и размеров матриц axpy_gold.cpp computeGold axpy.cu main randomInit printDiff runAxpy axpy_kernel.cu axpy (kernel)

void runAxpy(int argc, char** argv) инициализируем устройство (device) CUT_DEVICE_INIT(argc, argv); выделяем память на хосте для хранения векторов x и y unsigned int mem_size = sizeof(float) * n; float* h_x = (float*)malloc(mem_size); float* h_y = (float*)malloc(mem_size);

инициализируем векторы x и y случайными значениями randomInit(h_x, n); randomInit(h_y, n); выделяем память под векторы x и y на устройстве, копируем данные с хоста на устройство float* d_x; CUDA_SAFE_CALL(cudaMalloc((void**)&d_x, mem_size)); float* d_y; CUDA_SAFE_CALL(cudaMalloc((void**)&d_y, mem_size)); CUDA_SAFE_CALL(cudaMemcpy(d_x, h_x, mem_size, cudaMemcpyHostToDevice) ); CUDA_SAFE_CALL(cudaMemcpy(d_y, h_y, mem_size, cudaMemcpyHostToDevice) );

создаем и инициализируем таймер unsigned int timer = 0; CUT_SAFE_CALL(cutCreateTimer(&timer)); CUT_SAFE_CALL(cutStartTimer(timer)); определяем конфигурацию выполнения ядра (размер решетки блоков потоков и блока потоков) dim3 threads(BLOCK_SIZE); dim3 grid(n / threads.x); запускаем ядро копируем вычисленный вектор y с устройства на хост

останавливаем таймер, выводим время вычислений, освобождаем ресурсы таймера CUT_SAFE_CALL(cutStopTimer(timer)); printf("Processing time: %f (ms) \n", cutGetTimerValue(timer)); CUT_SAFE_CALL(cutDeleteTimer(timer)); вычисляем то же самое произведение на CPU float* reference = (float*) malloc(mem_size); computeGold(reference, alpha, h_x, h_y, n);

сравниваем результат, полученный на GPU, с результатом, полученным на CPU (по евклидовой норме) CUTBoolean res = cutCompareL2fe(reference, h_y, n, 1e-6f); printf("Test %s \n", (1 == res) ? "PASSED" : "FAILED"); if (res!=1) printDiff(reference, h_y, n); освобождаем память

__global__ void axpy( int n, float alpha, float* x, float* y) вычисляем координату текущего блока потоков и сохраняем её в переменную bid int bid = blockIdx.x вычисляем координату текущего потока в блоке потоков и сохраняем её в переменную tid вычисляем индекс элемента в исходном массиве, который будет обрабатываться текущим потоком int index = bid * BLOCK_SIZE + tid

вычисляем значение элемента массива, обрабатываемого текущим потоком y[index] = alpha * x[index] + y[index]

Nvidia CUDA Programming Guide Многочисленные курсы по CUDA: (на русском языке)

?