ГОСТ 28147-89 (Режим простой замены). (C++) - CUDA

В программе реализовывается алгоритм шифрования ГОСТ 28147-89 в режиме простой замены. Изначально хотел сделать шифрование и дешифрование на CPU и на GPU. На CPU проблем не возникло, работает в обе стороны без нареканий. За красоту кода и оптимизацию не ручаюсь, на моём уровне самое главное, чтобы программа просто работала. На GPU возникли сложности. Код компилится, все стадии проходит. Но на выходе, при шифровании, имеются результаты, отличные от результатов на CPU и расшифровать этот результат так же не получается. Т.е. работает не правильно.

Мб кто-то подскажет, по какой причине это может быть?

в файл "in.txt" для начальной проверки задаю ряд чисел "12345", чтобы их размер не превышал размерности одного блока.

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

Так же, сложность в том, что нет возможности установить NVIDIA Parallel Nsight и посмотреть, что там внутри происходит. Для файлов ввода вывода кодировка Windows 1251.

#define _CRT_SECURE_NO_WARNINGS
#include <cstdlib>
#include <stdio.h>
#include <conio.h>
#include <iostream>
#include <fstream>
#include <iomanip>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <cuda.h>
#include <device_functions.h>
#include <cuda_runtime_api.h>
#include <stdio.h>
#include <omp.h>
#define TH 1
unsigned char p[8][16] =   // таблица замен в памяти хоста
{
    { 0xc, 0x4, 0x6, 0x2, 0xa, 0x5, 0xb, 0x9, 0xe, 0x8, 0xd, 0x7, 0x0, 0x3, 0xf, 0x1 },
    { 0x6, 0x8, 0x2, 0x3, 0x9, 0xa, 0x5, 0xc, 0x1, 0xe, 0x4, 0x7, 0xb, 0xd, 0x0, 0xf },
    { 0xb, 0x3, 0x5, 0x8, 0x2, 0xf, 0xa, 0xd, 0xe, 0x1, 0x7, 0x4, 0xc, 0x9, 0x6, 0x0 },
    { 0xc, 0x8, 0x2, 0x1, 0xd, 0x4, 0xf, 0x6, 0x7, 0x0, 0xa, 0x5, 0x3, 0xe, 0x9, 0xb },
    { 0x7, 0xf, 0x5, 0xa, 0x8, 0x1, 0x6, 0xd, 0x0, 0x9, 0x3, 0xe, 0xb, 0x4, 0x2, 0xc },
    { 0x5, 0xd, 0xf, 0x6, 0x9, 0x2, 0xc, 0xa, 0xb, 0x7, 0x8, 0x1, 0x4, 0x3, 0xe, 0x0 },
    { 0x8, 0xe, 0x2, 0x5, 0x6, 0x9, 0x1, 0xc, 0xf, 0x4, 0xb, 0x0, 0xd, 0xa, 0x3, 0x7 },
    { 0x1, 0x7, 0xe, 0xd, 0x0, 0x5, 0x8, 0x3, 0x4, 0xf, 0xa, 0x6, 0x9, 0xc, 0xb, 0x2 }
};
// ключик в памяти хоста
unsigned __int32 key[8] =
{
    0x0123,
    0x4567,
    0x89AB,
    0xCDEF,
    0x0123,
    0x4567,
    0x89AB,
    0xCDEF
};
#pragma once
using std::cout;
using std::cin;
using std::endl;
//---------------------------------------------------------------------------
// взято из хелпа, определяем размер файла
long filesize(FILE *stream)
{
    long curpos, length;
    curpos = ftell(stream);
    fseek(stream, 0L, SEEK_END);
    length = ftell(stream);
    fseek(stream, curpos, SEEK_SET);
    return length;
}
__device__ unsigned char pD[8][16] =  // таблица замен в памяти устройства
{
{ 0xc, 0x4, 0x6, 0x2, 0xa, 0x5, 0xb, 0x9, 0xe, 0x8, 0xd, 0x7, 0x0, 0x3, 0xf, 0x1 },
{ 0x6, 0x8, 0x2, 0x3, 0x9, 0xa, 0x5, 0xc, 0x1, 0xe, 0x4, 0x7, 0xb, 0xd, 0x0, 0xf },
{ 0xb, 0x3, 0x5, 0x8, 0x2, 0xf, 0xa, 0xd, 0xe, 0x1, 0x7, 0x4, 0xc, 0x9, 0x6, 0x0 },
{ 0xc, 0x8, 0x2, 0x1, 0xd, 0x4, 0xf, 0x6, 0x7, 0x0, 0xa, 0x5, 0x3, 0xe, 0x9, 0xb },
{ 0x7, 0xf, 0x5, 0xa, 0x8, 0x1, 0x6, 0xd, 0x0, 0x9, 0x3, 0xe, 0xb, 0x4, 0x2, 0xc },
{ 0x5, 0xd, 0xf, 0x6, 0x9, 0x2, 0xc, 0xa, 0xb, 0x7, 0x8, 0x1, 0x4, 0x3, 0xe, 0x0 },
{ 0x8, 0xe, 0x2, 0x5, 0x6, 0x9, 0x1, 0xc, 0xf, 0x4, 0xb, 0x0, 0xd, 0xa, 0x3, 0x7 },
{ 0x1, 0x7, 0xe, 0xd, 0x0, 0x5, 0x8, 0x3, 0x4, 0xf, 0xa, 0x6, 0x9, 0xc, 0xb, 0x2 }
};
__device__ unsigned __int32 keyD[8] =   // ключ в памяти устройства
{
0x0123,
0x4567,
0x89AB,
0xCDEF,
0x0123,
0x4567,
0x89AB,
0xCDEF
};
__global__ void rpz_DEVICE(unsigned __int32* n1_d, unsigned __int32* n2_d, int* rezhim_d, int* block_D)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x; 
    unsigned __int32 SUM232 = 0;
    if (i < *block_D)
    {
        // 32 цикла простой замены
        // ключ считываем в требуемом ГОСТом порядке
        int c = 0;
        for (int k = 0; k<32; k++)
        {
            if (*rezhim_d == 1) { if (k == 24) c = 7; }
            else { if (k == 8) c = 7; }
            // суммируем в сумматоре СМ1
            SUM232 = keyD[c] + n1_d[i];
            // заменяем по таблице замен
            unsigned char first_byte = 0, second_byte = 0, zam_symbol = 0;
            int n = 7;
            for (int q = 3; q >= 0; q--)
            {
                zam_symbol = *((unsigned char *)&SUM232 + q);
                first_byte = (zam_symbol & 0xF0) >> 4;
                second_byte = (zam_symbol & 0x0F);
                first_byte = pD[n][first_byte];
                n--;
                second_byte = pD[n][second_byte];
                n--;
                zam_symbol = (first_byte << 4) | second_byte;
                *((unsigned char *)&SUM232 + q) = zam_symbol;
            }
            SUM232 = (SUM232 << 11) | (SUM232 >> 21); // циклический сдвиг на 11
            SUM232 = n2_d[i] ^ SUM232; // складываем в сумматоре СМ2
            if (k<31)
            {
                n2_d[i] = n1_d[i];
                n1_d[i] = SUM232;
            }
            if (*rezhim_d == 1)
            {
                if (k<24)
                {
                    c++;
                    if (c>7) c = 0;
                }
                else
                {
                    c--;
                    if (c<0) c = 7;
                }
            }
            else
            {
                if (k<8)
                {
                    c++;
                    if (c>7) c = 0;
                }
                else
                {
                    c--;
                    if (c<0) c = 7;
                }
            }
        }
        n2_d[i] = SUM232;
    }
}
// функция, реализующая работу ГОСТ 28147-89 в режиме простой замены
void rpz_HOST(int block_H, unsigned __int32 *n1, unsigned __int32 *n2, int rezh)
{
    unsigned __int32 SUM232 = 0;
    for (int i = 0; i<block_H; i++)
    {
        // 32 цикла простой замены
        // ключ считываем в требуемом ГОСТом порядке
        int c = 0;
        for (int k = 0; k<32; k++)
        {
            if (rezh == 1) { if (k == 24) c = 7; }
            else { if (k == 8) c = 7; }
            // суммируем в сумматоре СМ1
            SUM232 = key[c] + n1[i];
            // заменяем по таблице замен
            unsigned char first_byte = 0, second_byte = 0, zam_symbol = 0;
            int n = 7;
            for (int q = 3; q >= 0; q--)
            {
                zam_symbol = *((unsigned char *)&SUM232 + q);
                first_byte = (zam_symbol & 0xF0) >> 4;
                second_byte = (zam_symbol & 0x0F);
                first_byte = p[n][first_byte];
                n--;
                second_byte = p[n][second_byte];
                n--;
                zam_symbol = (first_byte << 4) | second_byte;
                *((unsigned char *)&SUM232 + q) = zam_symbol;
            }
            SUM232 = (SUM232 << 11) | (SUM232 >> 21); // циклический сдвиг на 11
            SUM232 = n2[i] ^ SUM232; // складываем в сумматоре СМ2
            if (k<31)
            {
                n2[i] = n1[i];
                n1[i] = SUM232;
            }
            if (rezh == 1)
            {
                if (k<24)
                {
                    c++;
                    if (c>7) c = 0;
                }
                else
                {
                    c--;
                    if (c<0) c = 7;
                }
            }
            else
            {
                if (k<8)
                {
                    c++;
                    if (c>7) c = 0;
                }
                else
                {
                    c--;
                    if (c<0) c = 7;
                }
            }
        }
        n2[i] = SUM232;
    }
}
//---------------------------------------------------------------------------
int main()
{
    // выбираем шифрование или расшифрование
    int HorD = 0;
    int rezhim = 0;
    setlocale(LC_ALL, "Russian");
    //Переменная которая будет хранить время выполнения программы на устройстве
    double start_execution_time = 0.0;
    double end_execution_time = 0.0;
    do
    {
        printf("Выберите режим работы:\nНа хосте      - 1\nНа устройстве - 2\n");
        scanf("%d", &HorD);
    } while ((HorD != 1) && (HorD != 2)); // повторяем до тех пор, пока не будет введено 1 или 
    do
    {
        printf("Выберите режим работы:\nШифрование    - 1\nРасшифрование - 2\n");
        scanf("%d", &rezhim);
    } while ((rezhim != 1) && (rezhim != 2)); // повторяем до тех пор, пока не будет введено 1 или 
    FILE *f_begin, *f_end; // потоки для исходного и конечного файлов
    char open_str[8] = "in.txt";
    char save_str[9] = "out.txt";
    char N_H[4]; // 32-разрядный накопитель,
    unsigned __int32 *n1, *n2; // накопители N1, N2, и сумматор
    // открываем файлы
    f_begin = fopen(open_str, "rb"); // Открывает двоичный файл для чтения
    f_end = fopen(save_str, "ab"); //дописывает информацию в двоичныйы файл для записи
    // определим количество блоков
    float blokoff;
    blokoff = 8 * filesize(f_begin);
    blokoff = blokoff / 64;
    int block_H = blokoff;
    if (blokoff - block_H>0) block_H++;
    int sh;
    if (filesize(f_begin) >= 4) sh = 4; else sh = filesize(f_begin);
    int sh1 = 0;
    int flag = 0;
    n1 = new unsigned __int32[block_H];
    n2 = new unsigned __int32[block_H];
    // начнем считывание и преобразование блоков
    // присутствуют проверки на полноту блоков, чтобы считать только нужное количество бит
    for (int i = 0; i<block_H; i++)
    {
        // записываем в накопитель N1
        for (int q = 0; q<4; q++) *((unsigned char *)&N_H + q) = 0x00;
        if ((sh1 + sh)<filesize(f_begin))
        {
            fread(N_H, sh, 1, f_begin);
            sh1 += sh;
        }
        else
        {
            sh = filesize(f_begin) - sh1;
            fread(N_H, sh, 1, f_begin);
            flag = 1;
        }
        n1[i] = *((unsigned __int32 *)&N_H);
        cout << endl;
        cout << "i = " << i << endl;
        cout << "n1[i] = " << n1[i] << endl;
        // записываем в накопитель N2
        for (int q = 0; q<4; q++) *((unsigned char *)&N_H + q) = 0x00;
        if ((sh1 + sh)<filesize(f_begin))
        {
            fread(N_H, sh, 1, f_begin);
            sh1 += sh;
        }
        else
        {
            if (flag == 0)
            {
                sh = filesize(f_begin) - sh1;
                fread(N_H, sh, 1, f_begin);
            }
        }
        n2[i] = *((unsigned __int32 *)&N_H);
        cout << "n2[i] = " << n2[i] << endl;
    }
    start_execution_time = omp_get_wtime(); // Начальная отсечка времени
    if (HorD == 1)
    {
        rpz_HOST(block_H, n1, n2, rezhim); // запускаем РПЗ
    }
    else
    {
        int blocks, thread;
        thread = TH;
        blocks = (block_H*block_H) / TH + 1;
        int* block_D;
        int *rezhim_d;
        unsigned __int32 *n1_d, *n2_d;  
        unsigned __int32 *n1_h, *n2_h;  
        int NumBytes = block_H*sizeof(unsigned __int32);
        cudaMallocHost((void**) &n1_h,      NumBytes    );
        cudaMallocHost((void**) &n2_h,      NumBytes    );
        cudaMalloc((void**) &n1_d,      NumBytes    );
        cudaMalloc((void**) &n2_d,      NumBytes    );
        cudaMalloc((void**) &rezhim_d,  sizeof(int) );
        cudaMalloc((void**) &block_D,   sizeof(int) );
        cudaMemcpy(n1_d, &n1, NumBytes, cudaMemcpyHostToDevice);
        cudaMemcpy(n2_d, &n2, NumBytes, cudaMemcpyHostToDevice);
        cudaMemcpy(block_D,     &block_H,   sizeof(int),    cudaMemcpyHostToDevice);
        cudaMemcpy(rezhim_d,    &rezhim,    sizeof(int),    cudaMemcpyHostToDevice);
        rpz_DEVICE << <blocks, thread >> >( n1_d,   n2_d,   rezhim_d,   block_D );
        cudaMemcpy(n1_h, n1_d, NumBytes, cudaMemcpyDeviceToHost);
        cudaMemcpy(n2_h, n2_d, NumBytes, cudaMemcpyDeviceToHost);
        for (int i = 0; i<block_H; i++)
        {
            n1[i] = n1_h[i];
            n2[i] = n2_h[i];
        }
    }
    end_execution_time = omp_get_wtime(); //Конечная отсечка времени
    printf("Время выполнения на программы:  %lf \n", end_execution_time - start_execution_time);
    for (int i = 0; i<block_H; i++)
    {
        cout << endl;
        cout << "i rez  = " << i << endl;
        cout << "n1[i] rez = " << n1[i] << endl;
        cout << "n2[i] rez = " << n2[i] << endl;
        // вывод результата в файл
        char sym_rez;
        for (int q = 0; q <= 3; q++)
        {
            sym_rez = *((unsigned char *)&n1[i] + q);
            fprintf(f_end, "%c", sym_rez);
        }
        for (int q = 0; q <= 3; q++)
        {
            sym_rez = *((unsigned char *)&n2[i] + q);
            fprintf(f_end, "%c", sym_rez);
        }
    }
    fclose(f_begin);
    fclose(f_end);
    system("pause");
    return 0;
}
//---------------------------------------------------------------------------