Вычисление НОД Nvidia CUDA

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "device_functions.h"
#include <stdio.h> //стандартный заголовочный файл ввода-вывода
#include <cstdlib> //заголовочный файл определяет функции динамического управления памятью, генерации случайных чисел, поиска, сортировки и преобразования типов данных
#include <iostream> //для вывода на экран
#include <string>   //для поддержки строк
#include <fstream> //для чтения файлов
 
using namespace std;
 
 
//чтение элементов из файла и их запись в массив
int readFromFile(unsigned int * a)
{
    int countOfLines = 0;
    ifstream inputFile("test.txt");
    string str;
    while (!inputFile.eof())
    {
        getline(inputFile, str);
        unsigned long x = stoul(str);
        a[countOfLines] = (unsigned int)x;
        countOfLines++;
    }
    inputFile.close();
    return countOfLines;
}
 
//вывод элементов в файл 
void writeInFile(unsigned int * a, int len, float time)
{
    ofstream outputFile;
    outputFile.open("ParallelResult.txt");
    for (int i = 0; i < len - 1; i++)
    {
        outputFile << a[i];
        outputFile << "\n";
    }
    outputFile << a[len - 1];
    outputFile.close();
}
 
//подсчет количества элементов в файле
int countElementsFromFile()
{
    int countOfLines = 0;
    ifstream inputFile("test.txt");
    string str;
 
    while (!inputFile.eof())
    {
        std::getline(inputFile, str);
        countOfLines++;
    }
 
    std::cout << "Count of lines:" + countOfLines;
    inputFile.close();
    return countOfLines;
}
 
//вывод элементов массива на экран
void showArrayElements(unsigned int *a, int len)
{
    for (int i = 1; i <= len; i++)
    {
        printf("%u ", a[i - 1]);
        if (i % 32 == 0) printf("\n");
    }
    printf("\n");
}
 
unsigned int* generateCompareArray(int n)
{
    unsigned int * compareArray;
    int numOfCompare = (n*n - n) / 2;
    //одномерный массив с индексами сравниваемых значений
    compareArray = new unsigned int[numOfCompare * 2];
    int temp = 0;
    //заполнение массива индексами сравниваемых значений
    for (int i = 0; i < n; i++)
    {
        for (int j = i + 1; j < n; j++)
        {
            compareArray[temp] = i;
            temp++;
            compareArray[temp] = j;
            temp++;
        }
    }
    return compareArray;
}
 
//Функция параллельного деления значения на 2
__device__ void divide2_Parallel(unsigned int *number, int id)
{
    unsigned int temp;
    if (id % 32 != 0)
    {
        temp = number[id - 1];
        number[id] = number[id] >> 1;
        number[id] = number[id] | (temp << 31);
    }
    else
    {
        temp = 0;
        number[id] = number[id] >> 1;
        number[id] = number[id] | (temp << 31);
    }
    delete &temp;
}
 
//Функция параллельного умножения значения на 2
__device__ void mult2_Parallel(unsigned int *number, int id)
{
    unsigned int temp;
    if (id % 31 != 0)
    {
        temp = number[id + 1];
        number[id] = number[id] << 1;
        number[id] = number[id] | (temp >> 31);
    }
    else
    {
        number[id] = number[id] << 1;
    }
    delete &temp;
}
 
//Функция параллельного вычитания
__device__ void sub_Parallel(unsigned int *number1, unsigned int *number2, int id, unsigned int *flags)
{
    if (number2[id] > number1[id])
    {
        number1[id] = +4294967295 - number2[id] + 1;
        flags[id - 1] += 1;
    }
    else
    {
        number1[id] -= number2[id];
    }
    do
    {
        if (flags[id] > 0)
        {
            if (number1[id] == 0)
            {
                number1[id] = 4294967295;
                flags[id - 1] += 1;
            }
            else
            {
                number1[id] -= 1;
                flags[id] -= 1;
            }
        }
    } while (flags[id] != 0);
}
 
//Функция проверки пары значений на больше или равно
__device__ bool greaterEq_Parallel(int *number1, int *number2, int id)
{
    int count = 0;
    for (int i = 0; i < 32; i++)
    {
        if ((unsigned)number1[i] < (unsigned)number2[i]) return false;
        else if ((unsigned)number1[i] > (unsigned)number2[i]) return true;
        else count++;
    }
    if (count == 32) return true;
    else return false;
}
 
//Функция проверки на нечетность значения 
__device__ bool nechet_Parallel(int *number1)
{
    int temp = number1[31] % 2;
    if (temp == 0) return false;
    else return true;
 
}
 
//Функция проверки на равенство значения нулю 
__device__ bool zero_Parallel(int *number1)
{
    for (int i = 0; i < 32; i++)
        if (number1[i] != 0) return false;
    return true;
}
 
//Функция параллельного копирования одного массива в другой
__device__ void copyToFrom_Parallel(int *to, int *from, int id)
{
    to[id] = from[id];
}
 
//Основная функция
__global__ void incKernel(unsigned int a[], unsigned int b[], unsigned int c[], int n, int numOfThreads)
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    int shift;
 
    __shared__ int tempArray[32];
    __shared__ int number1[32];
    __shared__ int number2[32];
    __shared__ unsigned int flags[32];
    __shared__ int pos;
 
 
    __shared__ int count;
 
    //копирование данных 2 сравниваемых значений из глобальной памяти в регистровую  
    number1[idx] = a[b[idx / 32 * 2] + idx % 32];
    number2[idx] = a[b[idx / 32 * 2 + 1] * 32 + idx % 32];
 
 
    //ВЫЧИСЛЕНИЕ НОД   --- НАЧАЛО
 
 
    //если число1 = 0, то НОД = число2, иначе наоборот
    if (zero_Parallel(number1))
    {
        //копирование результатов
        for (int j = 0; j < 32; j++)    c[j] = number2[j];
        return;
    }
    else if (zero_Parallel(number2))
    {
        //копирование результатов
        for (int j = 0; j < 32; j++)    c[j] = number1[j];
        return;
    }
 
    //ОСНОВНАЯ ЧАСТЬ
 
    for (shift = 0; !nechet_Parallel(number1) & !nechet_Parallel(number2); ++shift)
    {
        divide2_Parallel((unsigned*)number1, idx);
        divide2_Parallel((unsigned*)number2, idx);
    }
 
    while (!nechet_Parallel(number1))
    {
        divide2_Parallel((unsigned*)number1, idx);
    }
 
 
    do
    {
        while (!nechet_Parallel(number2))
        {
            divide2_Parallel((unsigned*)number2, idx);
        }
 
        if (!greaterEq_Parallel(number1, number2, idx))
        {
            sub_Parallel((unsigned*)number2, (unsigned*)number1, idx, flags);
        }
        else
        {
 
            copyToFrom_Parallel(tempArray, number1, idx);
            sub_Parallel((unsigned*)tempArray, (unsigned*)number2, idx, flags);
            copyToFrom_Parallel(number1, number2, idx);
            copyToFrom_Parallel(number2, tempArray, idx);
        }
 
        divide2_Parallel((unsigned*)number2, idx);
 
    } while (!zero_Parallel(number2));
 
    while (shift != 0)
    {
        mult2_Parallel((unsigned*)number1, idx);
        shift--;
    }
 
    //ВЫЧИСЛЕНИЕ НОД   --- КОНЕЦ
 
 
    //копирование результатов
    for (int j = 0; j < 32; j++)
    {
        c[j] = number1[j];
    }
}
 
int main(int argc, char * argv[])
{
    int n = countElementsFromFile();
    printf("Num of input elements: %d\n", n);
    int numBytes = n * sizeof(int);
    unsigned int * inputArray = new unsigned int[n];
    unsigned int * b = new unsigned int[n / 32 * n / 32 - n / 32];
    unsigned int * resultArray = new unsigned int[n / 2];
 
    readFromFile(inputArray);
    showArrayElements(inputArray, n);
 
    //генерация массива с индексами сравниваемых "ключей"
    unsigned int * array = generateCompareArray(n);
    showArrayElements(array, n / 32 * n / 32 - n / 32);
 
    unsigned int * adev = NULL;
    unsigned int * bdev = NULL;
    unsigned int * cdev = NULL;
    cudaMalloc((void **)&adev, numBytes);   //массив сравниваемызх элементов
    cudaMalloc((void **)&bdev, (n / 32 * n / 32 - n / 32) * sizeof(int)); //массив с индексами сравниваемых элементов
    cudaMalloc((void **)&cdev, (n / 2) * sizeof(int));  //массив с результатами
 
    int numOfThreads = 32;
    dim3 threads = dim3(numOfThreads, 1);
    dim3 blocks = dim3(1, 1);
    cudaEvent_t start, stop;
    float gpuTime = 0.0f;
 
    cudaEventCreate(&start);
    cudaEventCreate(&stop);
 
    cudaEventRecord(start, 0);
    cudaMemcpy(adev, inputArray, numBytes, cudaMemcpyHostToDevice);
    cudaMemcpy(bdev, array, (n / 32 * n / 32 - n / 32) * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(cdev, resultArray, (n / 2) * sizeof(int), cudaMemcpyHostToDevice);
    incKernel << <blocks, threads >> > (adev, bdev, cdev, n, numOfThreads);
    cudaMemcpy(resultArray, cdev, (n / 2) * sizeof(int), cudaMemcpyDeviceToHost);
    cudaEventRecord(stop, 0);
 
    cudaEventSynchronize(stop);
    cudaEventElapsedTime(&gpuTime, start, stop);
 
    printf("time spent executing by the GPU: %d seconds\n", (int)gpuTime * 1000);
    printf("%f", gpuTime);
    printf("----------------------------------------\n");
 
    cudaEventDestroy(start);
    cudaEventDestroy(stop);
 
    //запись в результатов в файл
    writeInFile(resultArray, n / 2, gpuTime);
    printf("Data has been written in file...\n");
 
    //печать результирующего массива на экран
    showArrayElements(resultArray, n / 2);
 
    cudaFree(adev);
    cudaFree(bdev);
    cudaFree(cdev);
 
    delete inputArray;
    delete b;
    delete resultArray;
 
    system("pause");
    return 0;
 
}