ее. Лаб 7. Гистограмма текста. Гистограмма текста

• Название: Гистограмма текста
• Дата: 23.04.2022
• Размер: 19.78 Kb.
• Имя файла: Лаб 7. Гистограмма текста.docx
• Тип: Инструкция; #492148

Гистограмма текста

GPU Teaching Kit – Accelerated Computing

Цели

Цель данной лабораторной – реализовать эффективный алгоритм вычисления гистограммы массива символов ASCII. Имеется 128 ASCII символов и каждый соответствует своему собственному интервалу. Счетчики интервалов представлены 32-битными беззнаковыми целыми без насыщения. Используйте подход, выполняющий приватизацию в разделяемой памяти для каждого блока нитей, а после – автоматически модифицирующий глобальную гистограмму.

Ход выполнения

Отредактируйте код, чтобы выполнить следующие действия:

• 
  выделить память на устройстве
  
• 
  скопировать память хоста на устройство
  
• 
  инициализировать размерности блока нитей и ядра
  
• 
  запустить ядро CUDA
  
• 
  скопировать результат на хост
  
• 
  освободить память устройства
  

Места, куда следует вставить код, помечены комментариями //@@.

Инструкция по установке

Последнюю версию исходного кода лабораторной, наряду со скриптами сборки, можно найти в репозитории Bitbucket. Инструкции к Cmake и сборки лабораторной можно найти в файле README в корневой директории репозитория.

Исполняемый файл, являющийся результатом компиляции лабораторной, можно запустить следующей командой:

./Histogram_Template -e <expected.raw> \
-i <input.txt> -o <output.raw> -t integral_vector

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

Вопросы

1. 
   Опишите все оптимизации, которые вы попробовали выполнить, вне зависимости от того, стали ли вы их использовать в дальнейшем или нет, и вне зависимости от того, как они повлияли на производительность.
   
2. 
   С какими сложностями вы столкнулись во время выполнения оптимизации?
   
3. 
   Какая оптимизация дала наибольший прирост производительности?
   
4. 
   Сколько чтений из глобальной памяти выполнит выше ядро вычисления гистограммы? Объясните.
   
5. 
   Сколько записей в глобальную память выполнит выше ядро вычисления гистограммы? Объясните.
   
6. 
   Сколько атомарных операций выполнит выше ядро вычисления гистограммы? Объясните.
   
7. 
   Большинство текстов содержит только буквы, цифры и пробелы. Эти 95 знаков в таблице ASCII имеют последовательные значения кода: 32 - 126. Что можно сказать о ситуации (порядок выполнения) с атомарным доступом, когда более 95 нитей одновременно пытаются инкрементировать одну гистограмму?
   

Шаблон кода

Представленный код предлагается как отправная точка. Импорты, экспорты и проверка решения уже представлены в коде. Требуется вставить свой код в области, обозначенные //@@. Остальной код трогать не нужно. Руководство описывает функционал методов с префиксом wb*.

#include

#define NUM_BINS 4096

#define CUDA_CHECK(ans) \
{ gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, const char *file, int line,
bool abort = true) {
if (code != cudaSuccess) {
fprintf(stderr, "GPUassert: %s %s %d\n", cudaGetErrorString(code),
file, line);
if (abort)
exit(code);
}
}

int main(int argc, char *argv[]) {
wbArg_t args;
int inputLength;
unsigned int *hostInput;
unsigned int *hostBins;
unsigned int *deviceInput;
unsigned int *deviceBins;

args = wbArg_read(argc, argv);

wbTime_start(Generic, "Importing data and creating memory on host");
hostInput = (unsigned int *)wbImport(wbArg_getInputFile(args, 0),
&inputLength, "Integer");
hostBins = (unsigned int *)malloc(NUM_BINS * sizeof(unsigned int));
wbTime_stop(Generic, "Importing data and creating memory on host");

wbLog(TRACE, "The input length is ", inputLength);
wbLog(TRACE, "The number of bins is ", NUM_BINS);

wbTime_start(GPU, "Allocating GPU memory.");
//@@ Выделите память GPU
CUDA_CHECK(cudaDeviceSynchronize());
wbTime_stop(GPU, "Allocating GPU memory.");

wbTime_start(GPU, "Copying input memory to the GPU.");
//@@ Скопируйте память с хоста на GPU
CUDA_CHECK(cudaDeviceSynchronize());
wbTime_stop(GPU, "Copying input memory to the GPU.");

// Запуск ядра
// ----------------------------------------------------------
wbLog(TRACE, "Launching kernel");
wbTime_start(Compute, "Performing CUDA computation");
//@@ Выполните вычисления в ядре
wbTime_stop(Compute, "Performing CUDA computation");

wbTime_start(Copy, "Copying output memory to the CPU");
//@@ Скопируйте память обратно с GPU на хост
CUDA_CHECK(cudaDeviceSynchronize());
wbTime_stop(Copy, "Copying output memory to the CPU");

wbTime_start(GPU, "Freeing GPU Memory");
//@@ Освободите память GPU
wbTime_stop(GPU, "Freeing GPU Memory");

// Проверка корректности результатов
// -----------------------------------------------------
wbSolution(args, hostBins, NUM_BINS);

free(hostBins);
free(hostInput);
return 0;
}