Студопедия

КАТЕГОРИИ:

АстрономияБиологияГеографияДругие языкиДругоеИнформатикаИсторияКультураЛитератураЛогикаМатематикаМедицинаМеханикаОбразованиеОхрана трудаПедагогикаПолитикаПравоПсихологияРиторикаСоциологияСпортСтроительствоТехнологияФизикаФилософияФинансыХимияЧерчениеЭкологияЭкономикаЭлектроника


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




Ниже, с комментариями, приведён текст той части программы, которая предназначена для центрального процессора. Эта программа содержит операторы, специфические для 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

unsigned int mask_b; // параметр b

unsigned int mask_c; // параметр c

unsigned int seed; /* Константа, используемая для генерации инициализирующей последовательности {x0, ¼, xn-1} */

} mt_parameters;

 

// Параметры, одинаковые для всех параллельных генераторов

#define MT_MM 9 // Параметр m

#define MT_NN 19 // Параметр n

#define MT_WMASK 0xFFFFFFFFU /* «Маска для получения первых 32 битов числа */

#define MT_UMASK 0xFFFFFFFEU /* «Маска для получения 31 бита числа */

#define MT_LMASK 0x1U /* «Маска для получения 1 последнего бита числа */

#define MT_SHIFT0 12 // Параметр u

#define MT_SHIFTB 7 // Параметр s

#define MT_SHIFTC 15 // Параметр t

#define MT_SHIFT1 18 // Параметр l

 

#endif

#endif

 

На центральном процессоре исполняются также процедуры loadMTGPU и seedMTGPU, первая из которых загружает файл с инициализирующими параметрами для всех параллельных генераторов случайных чисел, а вторая – осуществляет инициализацию генераторов. Мы выделили эти процедуры в отдельный файл “MТ_Initialization.cu”, текст которого приведён ниже.

 

/* Текст файла “MТ_Initialization.cu” */

 

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

__device__ static mt_parameters device_MT[Number_of_Threads];

static mt_parameters host_MT[Number_of_Threads];

 

/* Процедура, загружающая параметры для инициализации генераторов случайных чисел из файла fname */

void loadMTGPU(const char *fname)

{

/* Открытие файла fname */

FILE *fd = fopen(fname, "rb");

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

if(!fd){

printf("initMTGPU(): failed to open %s\n", fname);

exit(0);

}

/* Загрузка содержимого файла в массив host_MT. Вывод сообщения и выход из программы в случае ошибки */

if( !fread(host_MT, sizeof(host_MT), 1, fd) ){

printf("initMTGPU(): failed to load %s\n", fname);

exit(0);

}

 

/* закрытие файла fname */

fclose(fd);

}

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

void seedMTGPU(unsigned int seed)

{

int i;

/* В оперативной памяти компьютера создаётся временный массив для хранения инициализирующих параметров - структур пользовательского типа mt_parameters */

mt_parameters *MT = (mt_parameters *)malloc(Number_of_Threads * sizeof(mt_parameters));

 

/* Цикл по всем Number_of_Threads = 4096 генераторам случайных чисел. Здесь они обрабатываются последовательно, так как процедура исполняется центральным процессором. Инициализирующий массив host_MT[i], загруженный из файла процедурой loadMTGPU (и модифицированный переменной seed) копируется в область памяти, на которую указывает указатель MT */

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

{

MT[i] = host_MT[i];

MT[i].seed = seed;

}

 

/* Область оперативной памяти, заданная указателем MT и содержащая инициализирующий массив, копируется в область видеопамяти (определяемую указателем device_MT), где она будет доступна графическому процессору */

CUDA_SAFE_CALL( cudaMemcpyToSymbol(device_MT, MT, sizeof(host_MT)) );

 

/* Освобождение оперативной памяти, которую занимал массив MT */

free(MT);

}


Поделиться:

Дата добавления: 2015-09-13; просмотров: 70; Мы поможем в написании вашей работы!; Нарушение авторских прав





lektsii.com - Лекции.Ком - 2014-2024 год. (0.006 сек.) Все материалы представленные на сайте исключительно с целью ознакомления читателями и не преследуют коммерческих целей или нарушение авторских прав
Главная страница Случайная страница Контакты