РВС. 8 РВС. Лабораторная работа 8 Гистограмма Распределенные вычислительные системы Специальность 10. 03. 01 Информационная безопасность
Скачать 366.48 Kb.
|
Министерство цифрового развития, связи и массовых коммуникаций Российской Федерации Ордена Трудового Красного Знамени федеральное государственное бюджетное образовательное учреждение «Московский технический университет связи и информатики» Лабораторная работа №8 «Гистограмма» «Распределенные вычислительные системы» Специальность 10.03.01 – Информационная безопасность Выполнил: студент группы БИБ2003 Абдиев Н. Проверил: Соловьёв А.С. Москва 2021 Цель работы: реализовать эффективный алгоритм для вычисления гистограммы входного массива целых чисел из определенного диапазона Исходный код: # include #define NUM_BINS 127 #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); } } __global__ void histogram_kernel(unsigned int* input, unsigned int* bins, unsigned int num_elements, unsigned int num_bins) { unsigned int tid = blockIdx.x * blockDim.x + threadIdx.x; extern __shared__ unsigned int bins_s[]; for (unsigned int binIdx = threadIdx.x; binIdx < num_bins; binIdx += blockDim.x) { bins_s[binIdx] = 0; } __syncthreads(); // Histogram for (unsigned int i = tid; i < num_elements; i += blockDim.x * gridDim.x) { atomicAdd(&(bins_s[input[i]]), 1); } __syncthreads(); // Фиксация в глобальной памяти for (unsigned int binIdx = threadIdx.x; binIdx < num_bins; binIdx += blockDim.x) { atomicAdd(&(bins[binIdx]), bins_s[binIdx]); } } __global__ void convert_kernel(unsigned int* bins, unsigned int num_bins) { unsigned int tid = blockIdx.x * blockDim.x + threadIdx.x; if (tid < num_bins) { bins[tid] = min(bins[tid], 127); } } void histogram(unsigned int* input, unsigned int* bins, unsigned int num_elements, unsigned int num_bins) { CUDA_CHECK(cudaMemset(bins, 0, num_bins * sizeof(unsigned int))); // Запустите ядро гистограммы { dim3 blockDim(512), gridDim(30); histogram_kernel <<< gridDim, blockDim, num_bins * sizeof(unsigned int)>>> ( input, bins, num_elements, num_bins); CUDA_CHECK(cudaGetLastError()); CUDA_CHECK(cudaDeviceSynchronize()); } // Убедитесь, что значения ячеек не слишком велики { dim3 blockDim(512); dim3 gridDim((num_bins +blockDim.x - 1) / blockDim.x); convert_kernel <<< gridDim, blockDim >>> (bins, num_bins); CUDA_CHECK(cudaGetLastError()); CUDA_CHECK(cudaDeviceSynchronize()); } } 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."); //Выделять память графического процессора CUDA_CHECK(cudaMalloc((void**)&deviceInput, inputLength * sizeof(unsigned int))); CUDA_CHECK( cudaMalloc((void**)&deviceBins, NUM_BINS * sizeof(unsigned int))); CUDA_CHECK(cudaDeviceSynchronize()); wbTime_stop(GPU, "Allocating GPU memory."); wbTime_start(GPU, "Copying input memory to the GPU."); //Скопируйте память в графический процессор CUDA_CHECK(cudaMemcpy(deviceInput, hostInput, inputLength * sizeof(unsigned int), cudaMemcpyHostToDevice)); CUDA_CHECK(cudaDeviceSynchronize()); wbTime_stop(GPU, "Copying input memory to the GPU."); // Запустить ядро // ---------------------------------------------------------- wbLog(TRACE, "Launching kernel"); wbTime_start(Compute, "Performing CUDA computation"); histogram(deviceInput, deviceBins, inputLength, NUM_BINS); wbTime_stop(Compute, "Performing CUDA computation"); wbTime_start(Copy, "Copying output memory to the CPU"); //Скопируйте память графического процессора обратно в центральный процессор здесь CUDA_CHECK(cudaMemcpy(hostBins, deviceBins, NUM_BINS * sizeof(unsigned int), cudaMemcpyDeviceToHost)); CUDA_CHECK(cudaDeviceSynchronize()); wbTime_stop(Copy, "Copying output memory to the CPU"); wbTime_start(GPU, "Freeing GPU Memory"); // Освободите память графического процессора здесь CUDA_CHECK(cudaFree(deviceInput)); CUDA_CHECK(cudaFree(deviceBins)); wbTime_stop(GPU, "Freeing GPU Memory"); // ----------------------------------------------------- wbSolution(args, hostBins, NUM_BINS); free(hostBins); free(hostInput); return 0; } Тест программы: Рисунок 1 – Скриншот результата работы программы Рисунок 2 – Скриншот входных данных Рисунок 3 – Скриншот работы программы Вывод: в результате выполнения лабораторных работ, я реализовал эффективный алгоритм вычисления гистограммы массива целых чисел Контрольные вопросы: |