Текст программы для центрального процессора

Ниже, с комментариями, приведён текст той части программы, которая предназначена для центрального процессора. Эта программа содержит операторы, специфические для CUDA, необходимые для взаимодействия центрального и графического процессоров. Она компилируется компилятором CUDA.

Прототипом программы был пример реализации генератора случайных чисел Mersenne Twister, предложенный компанией NVIDIA в дистрибутиве CUDA [15]. Этот текст был модифицирован в соответствии с решаемой задачей о проницаемости пластины.

/* Подключение стандартных библиотек, необходимых компилятору CUDA. Осуществляется по аналогии с примером [15] из дистрибутива CUDA */

#include <stdlib.h>

#include <stdio.h>

#include <time.h>

#include <string.h>

#include <cutil.h>

/* Подключение файлов, которые содержат:

à значения общих параметров генератора Mersenne Twister (MersenneTwister.h);

à процедуры для инициализации параллельных генераторов случайных чисел (MТ_Initialization.cu);

à текст вычислительного ядра (Monte_Carlo_kernel.cu) */

#include "MersenneTwister.h"

#include "MТ_Initialization.cu"

#include "Monte_Carlo_kernel.cu"

/* Произвольная константа, участвующая в инициализации генераторов случайных чисел. В данном примере она одинакова для всех генераторов, но можно задавать её для каждого генератора случайным образом */

const unsigned int SEED = 777;

/* Установка количества вычислительных потоков, каждый из которых будет моделировать движение одного нейтрона «с весом» через пластину */

const int Number_of_Threads = 32 * 128; // = 4096

/* Определение объёма памяти, необходимого для сохранения результатов расчёта: 2 числа по 4 байта на каждый поток */

const int OutPut_Memory = 2 * 4 * Number_of_Threads;

/* Процедура, с которой начинается исполнение приложения */

int main(int argc, char **argv)

{

/* Описание указателей на массивы чисел типа float (вещественные числа одинарной точности), которые будут размещены в видеопамяти и в оперативной памяти компьютера для записи результатов моделирования */

float *d_Rand; // Указатель для видеопамяти

float *h_RandGPU // Указатель для оперативной памяти;

/* Инициализация графического процессора */

CUT_DEVICE_INIT(argc, argv);

/* Выделение области для хранения результатов в оперативной памяти компьютера. h_RandGPU – указатель на начало этой области, эквивалентный названию массива. Тип (float *) означает, что в массиве будут храниться вещественные числа одинарной (32-битной) точности. */

h_RandGPU = (float *)malloc(OutPut_Memory);

/* Выделение такой же области для результатов в видеопамяти графического процессора. На начало этой области будет указывать указатель d_Rand. */

CUDA_SAFE_CALL(cudaMalloc((void **)&d_Rand, OutPut_Memory));

/* Автоматическое определение полного пути к файлу MersenneTwister.dat, в котором хранятся инициализирующие параметры для 4096 параллельных генераторов случайных чисел */

const char *dat_path = cutFindFilePath("MersenneTwister.dat", argv[0]);

/* Загрузка этого файла и выполнение процедуры, создающей инициализирующие последовательности параметров для всех 4096 генераторов случайных чисел */

loadMTGPU(dat_path);

seedMTGPU(SEED);

/* Синхронизация вычислительных потоков перед запуском графического процессора на исполнение вычислительного ядра */

CUDA_SAFE_CALL( cudaThreadSynchronize() );

/* Запуск вычислительного ядра на GPU. Сам текст процедуры, названной GPU_Monte_Carlo, находится в отдельном файле "Monte_Carlo_kernel.cu", который будет рассмотрен ниже, в п. 3.2.3. Инструкция <<<32, 128>>> означает, что будет запущено 32 «связки», по 128 потоков в каждой. Если программа будет исполняться на процессоре G 80 или предыдущих, то на каждый мультипроцессор придётся не менее двух «связок». Количество потоков 128 кратно 64, что является одним из условий оптимизации вычислений (см. п. 3.1.3).*/

GPU_Monte_Carlo<<<32, 128>>>(d_Rand);

/* Проверка успешности исполнения вычислительного ядра на GPU и вывод на экран сообщения в случае ошибки */

CUT_CHECK_ERROR("RandomGPU() execution failed\n");

/* Синхронизация вычислительных потоков после завершения расчёта на графическом процессоре */

CUDA_SAFE_CALL( cudaThreadSynchronize() );

/* Копирование результатов моделирования из видеопамяти в оперативную память компьютера. Область видеопамяти, начало которой показывает указатель d_Rand, копируется в область оперативной памяти, на которую указывает указатель h_RandGPU */

CUDA_SAFE_CALL( cudaMemcpy(h_RandGPU, d_Rand, OutPut_Memory, cudaMemcpyDeviceToHost) );

/* Обработка результатов моделирования по формулам (2.14)-(2.17). Исполняется без распараллеливания, на центральном процессоре */

float Reflected_weight = 0.0f; // Суммарный вес нейтронов, отразившихся
// от пластины

float Penetrating_weight = 0.0f; // Суммарный вес нейтронов, прошедших
// через пластину

int MC_address; // Адрес очередного нейтрона в массиве

for(int i = 0; i < Number_of_Threads; i++)

{

MC_address = 2 * i;

if (h_RandGPU[MC_address] < 0.5f) {

Reflected_weight += h_RandGPU[MC_address+1];

}

if (h_RandGPU[MC_address] > 0.5f) {

Penetrating_weight += h_RandGPU[MC_address+1];

}

} // Завершение анализа массива результатов

/* Расчёт и вывод на экран вероятностей прохождения пластины, отражения от пластины и поглощения в пластине, с учётом формул - */

printf("Probability of penetration: %E \n", Penetrating_weight / (float)Number_of_Threads);

printf("Probability of reflection : %E \n", Reflected_weight / (float)Number_of_Threads);

printf("Probability of absorption : %E \n", 1.0f – (Penetrating_weight + Reflected_weight) / (float)Number_of_Threads);

printf("Probability of deflection : %E \n", (MC_weight_overall - MC_weight_through) / (float)MC_out_neutrons);

/* Освобождение видеопамяти и оперативной памяти компьютера */

CUDA_SAFE_CALL( cudaFree(d_Rand) );

free(h_RandGPU);

/* Завершение работы программы */

CUT_EXIT(argc, argv);

}

Описания пользовательских типов, а также константы, постоянного доступа к которым программисту не требуется, часто хранят в так называемых заголовочных файлах, с традиционным расширением «.h». В нашем примере используется заголовочный файл “MersenneTwister.h”, где описана структура mt_parameters, используемая для инициализации параллельных генераторов MersenneTwister, а также задаются значения постоянных для этого генератора. Текст файла “MersenneTwister.h” приведён ниже.

/* Заголовочный файл “MersenneTwister.h” */

#ifndef MERSENNETWISTER_H

#define MERSENNETWISTER_H

#ifndef mersennetwister_h

#define mersennetwister_h

#define DCMT_SEED 4172

#define MT_RNG_PERIOD 607

/* Определяемые ниже переменные фигурируют в коде программы вычислительного ядра (п. 3.2.3). Структура mt_parameters содержит варьируемые параметры, различные для каждого из параллельных генераторов случайных чисел */

typedef struct

{

unsigned int matrix_a; // параметр a