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

• Название: Инструкция по установке Последнюю версию исходного кода лабораторной, наряду со скриптами сборки, можно найти в репозитории Bitbucket. Инструкции к Cmake и сборки лабораторной можно найти в файле readme в корневой директории репозитория
• Анкор: лаб 6
• Дата: 23.04.2022
• Размер: 20.17 Kb.
• Имя файла: Лаб 6. Гистограмма.docx
• Тип: Инструкция; #492146

Гистограмма

GPU Teaching Kit – Accelerated Computing

Цели

Цель данной лабораторной – реализовать эффективный алгоритм для вычисления гистограммы входного массива целых чисел из определенного диапазона. Каждое целое будет засчитано в одном из столбцов гистограммы, следовательно, значения будут варьироваться от 0 до (NUM_BINS - 1). Счетчики интервалов должны быть представлены беззнаковыми 32-битными целыми и должны насыщаться на значении 127 (переставать увеличиваться по достижении этого числа). Длина входного массива будет меньше, чем 2^32. NUM_BINS равно 4096 в этой лабораторной.

Задача может быть разбита на два ядра, первое выполняет вычисление гистограммы без учета насыщения, второе подправляет значения на интервалах, если они слишком большие. Однако оба этих шага можно объединить в одно ядро.

Предварительные требования

Прежде чем начать, убедитесь, что:

• 
  вы изучили все лекции и материалы по Модулю 7.
  

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

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

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

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

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

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

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

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

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

Вопросы

1. 
   Опишите все оптимизации, которые вы попробовали выполнить, вне зависимости от того, стали ли вы их использовать в дальнейшем или нет, и вне зависимости от того, как они повлияли на производительность.
   

2. 
   С какими сложностями вы столкнулись во время выполнения оптимизации?
   

3. 
   Какая оптимизация дала наибольший прирост производительности?
   

4. 
   Сколько чтений из глобальной памяти выполнит ваше ядро вычисления гистограммы? Объясните.
   

5. 
   Сколько записей в глобальную память выполнит ваше ядро вычисления гистограммы? Объясните.
   

6. 
   Сколько атомарных операций выполнит выше ядро вычисления гистограммы? Объясните.
   

7. 
   Сколько конкурентных ожидается, если все элементы входного массива будут содержать одинаковые значения?
   

8. 
   Сколько конкурентных ожидается, если все элементы входного массива будут содержать случайные значения?
   

Шаблон кода

Представленный код предлагается как отправная точка. Импорты, экспорты и проверка решения уже представлены в коде. Требуется вставить свой код в области, обозначенные //@@. Остальной код трогать не нужно. Руководство описывает функционал методов с префиксом 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;
}