CUDA C - CRC32 - 寻找未知多项式和 crcxor - 程序加速

Posted 2023-04-15

【中文标题】CUDA C - CRC32 - 寻找未知多项式和 crcxor - 程序加速

【英文标题】：CUDA C - CRC32 - Finding unknown polynom and crcxor - program speed up

【发布时间】：2021-12-29 19:59:58

【问题描述】：

我一直在寻找与我的问题相关的问题，但只找到了有关 CRC32 反转的问题。我的主题有点不同。

我是一个新手程序员，我有这样的任务要做。我有输入（3 个 4 字节的字符串）。对于这些数据，我知道使用类似于 CRC32 的哈希函数计算的三个校验和。但是，它不是标准的 CRC32，因为它在多项式和 crcxor 参数的默认值和未知值之间存在差异。 因此，对于 4 字节的输入数据，我使用从 0 到 0xFFFFFFFF 的多项式的不同值以及使用介于 0 和 0xFFFF 之间的值的参数 crcxor 的不同值来计算 CRC。我用 CUDA C 编写了这个程序，因为它比 CPU 运行得更快。这是我在“Hello World”和“VectorAdd”之后的第三个 CUDA C 程序:)。要计算所有可能的 0xFFFF x 0xFFFFFFFF 变体，我的 NVIDIA GTX1060 卡大约需要 5 个小时。 我想问是否可以修改或优化以下程序代码以更快地完成这项任务？

最终，我想计算 0xFFFFFFFF x 0xFFFFFFFF 但我还不知道是否可以在短时间内完成。

如果有人想看看我的程序代码并提供有价值的反馈，我将非常感激。

 #include <stdio.h>
#include <cuda_runtime.h>
#include <cuda_profiler_api.h>


__device__ unsigned long calculate_crc(unsigned long data, unsigned long poly, unsigned long cxor) 
// truncated function for constant values crcinit = 0 refin = 0 refout = 0 direct = 0

    unsigned long i, j, k, c, bit, crc = 0;
    for (i=0,k=24; i<4; i++,k-=8) 
    
        c = (data>>k)&0xFF;
        for (j=0x80; j; j>>=1) 
        
            bit = crc & 0x80000000;
            crc<<= 1;
            if (c & j) bit^= 0x80000000;
            if (bit)   crc^= poly;
        
       
    crc^= cxor;
    crc&= 0xFFFFFFFF;   
    return crc;



__global__ void calculate_crc_parameters(unsigned long n) 

    unsigned long polynom = 0; 
    unsigned long crcxor = 0;
    
    //Input data:
    const unsigned long data1 = 0x928F640C;
    const unsigned long data2 = 0x0121B30E;
    const unsigned long data3 = 0xCB652607;
    
    // calculated CRC for the above input data and for polynom 0xFF7A1DB7: crcinit: 0, crcxor: 0x00000000, refin: 0, refout: 0, direct: 0:
    // for these CRCs, the function should find the polynomial 0xFF7A1DB7 and crcxor = 0
    // finds it right away because crcxor = 0
    const unsigned long crc1 = 0x7076BCEB;  
    const unsigned long crc2 = 0x1F719D7A;  
    const unsigned long crc3 = 0x8369D986;
    
    // other example crc - for crcxor> 0
    // computed CRC for polynom 0xFF7A1DB7: crcinit: 0, crcxor: 0x000000FF, refin: 0, refout: 0, direct: 0:
    // for these CRCs, the function should find the polynomial 0xFF7A1DB7 and crcxor = 0x000000FF
    // Program find it after 1m 12sec.
    /*
    const unsigned long crc1 = 0x7076BC14;  
    const unsigned long crc2 = 0x1F719D85;  
    const unsigned long crc3 = 0x8369D979;
    */  
    
    // computed CRC for polynom 0xFF7A1DB7: crcinit: 0, crcxor: 0x0000FFFE, refin: 0, refout: 0, direct: 0:
    // for these CRCs, the function should find the polynomial 0xFF7A1DB7 and crcxor = 0x0000FFFE
    // searches for 5 hours
    /*
    const unsigned long crc1 = 0x70764315;  
    const unsigned long crc2 = 0x1F716284;  
    const unsigned long crc3 = 0x83692678;
    */
    
    // CRCs - polynom 0xFF7A1DB7: crcinit: 0, crcxor: 0xFF7A1DB7, refin: 0, refout: 0, direct: 0:
    // no implementation for 8-byte crcxor yet - and it would count for a long time
    /*
    const unsigned long crc1 = 0x8F0CA15C;  
    const unsigned long crc2 = 0xE00B80CD;  
    const unsigned long crc3 = 0x7C13C431;
    */
    unsigned int index_x  = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int stride_x = blockDim.x * gridDim.x;
    unsigned int index_y  = blockIdx.y * blockDim.y + threadIdx.y;
    unsigned int stride_y = blockDim.y * gridDim.y;
    unsigned int index_z  = blockIdx.z * blockDim.z + threadIdx.z;
    unsigned int stride_z = blockDim.z * gridDim.z;
    
    if((index_x<n)&&(index_y<n)&&(index_z<n))
    
        polynom = (index_x << 16) ^ index_y; // "gluing" the polynomial
        // to get polynom e.g. 0xFF7A1DB7 we have to "glue it" with index_x and index_y
        // if index_x == 0xFF7A then LSH by 16 places and we get 0xFF7A0000
        // then xor from index_y: 0xFF7A0000 xor 0x00001DB7 and is 0xFF7A1DB7
        crcxor = index_z; // crcxor will take the values of index_z that is from 0x0000 to 0xFFFF
        if(calculate_crc(data1,polynom,crcxor)==crc1)
            if(calculate_crc(data2,polynom,crcxor)==crc2)
                if(calculate_crc(data3,polynom,crcxor)==crc3)       // compute three checksums and compare them
                    printf("\nCRC parameters found ---> polynom: 0x%08X,  crcxor: 0x%08X\n", polynom,crcxor);
                    // if the calculated 3 crc sums agree with the known 3 crcs, then display the parameters for which they were calculated
    
        if ((crcxor%0xFF==0)&&(polynom==0xFFFFFFFF)) printf("@");       // 1m 12s from displaying @ to the next @
        // if the @ sign is displayed 256 times, this will be the end of the program
        index_x+=stride_x;
        index_y+=stride_y;
        index_z+=stride_z;
    




int main(void)

    unsigned long N = 0x10000;          //  0xFFFF + 0x01 = 65536dec
    ////////////////////////////////////////////////
    // for computing only in X and Y axes - for crcxor = zero all the time
    dim3 dimBlock(   4,     4,    1);
    dim3 dimGrid(16384, 16384,    1);
    
    ////////////////////////////////////////////////
    // for computing on the X, Y and Z axes, i.e. for crcxor taking values from the Z axis from 0 to 65535
    //dim3 dimBlock(   4,     4,   64);             //  4 * 4 * 64 = 1024 --- maximum block size
    //dim3 dimGrid(16384, 16384, 1024);            //uncomment this 2 lines for crcxor > 0
    //               4      4    64
    //               *      *     *
    //           16384  16384  1024
    //               =      =     =
    //        0x10000 0x10000 0x10000
    //   x, y, and z will trigger 65,536 times each
    
    cudaProfilerStart();    
    calculate_crc_parameters<<<dimGrid, dimBlock>>>(N);
    cudaDeviceSynchronize();
    cudaDeviceReset();
    cudaProfilerStop();
    return 0;

我在 cmd 中编译：nvcc name.cu -o name

我使用 Cuda Toolkit 11.5 在 win10 上工作 显卡是 NVIDIA GTX 1060。

指针或内存分配的使用能以某种方式加速这个程序吗？

我计算测试CRC值here

【问题讨论】：

您可以从使用合理的块大小开始。通过每个块使用 16 个线程，您选择仅使用 GPU 理论计算能力的 50%。 CUDA 经线大小为 32。块大小应始终为经线大小的整数倍

所以像这样声明 dim 变量：dim3 dimBlock(32, 32,1); dim3 dimGrid(2048, 2048, 1); 比声明 dim3 dimBlock( 4, 4, 1); dim3 dimGrid(16384, 16384, 1); 更好？？？在这两种情况下，x 和 y 都是 65536，对吧？

是的，在每个块的线程数小于32的情况下

【参考方案1】：

优化应该从算法开始，而不是优化一种痛苦无意义的蛮力方法。

您可以考虑多项式和最终异或的搜索，首先执行多项式，然后（通常）找到异或值。您需要做的就是取两个数据值的异或，并找到产生这些值的两个 CRC 的异或的多项式，假设最终异或为零。您将需要尝试至少两对才能将其缩小到多项式的一种选择。

一旦你有了多项式，现在就计算你的一个数据值的 CRC，用所需的 CRC 异或，现在你有了最终的异或值。第二步无需搜索。

多项式搜索速度足够快，您只需使用 CPU。没有 GPU 或 CUDA 或任何需要的东西。在我用了三年的笔记本电脑上花了 40 秒。你只需要尝试奇数多项式。甚至多项式也是无效的。

异或数据和 CRC 也会取消初始值。因此，对于具有非零初始值和非零最终异或的 CRC，您可以通过这种方式找到多项式。然而，为了解决初始值和最终异或，您将需要具有不同长度消息的示例，即除了所有四字节消息之外。初始值和最终异或的 2³² 种可能组合将匹配四字节消息的任何和所有 CRC。

顺便说一句，您的 CRC 例程是不必要的复杂。请参见下面的等效项。这将打印poly = ff7a1db7, xor = 0000fffe:

#include <stdio.h>
#include <stdint.h>

uint32_t calculate_crc(uint32_t data, uint32_t poly, uint32_t xor) 
    for (int i = 0; i < 32; i++)
        data = data & 0x80000000 ? (data << 1) ^ poly : data << 1;
    return data ^ xor;


void findp(uint32_t data1, uint32_t data2, uint32_t data3,
           uint32_t crc1, uint32_t crc2, uint32_t crc3) 
    uint32_t d = data2, c = crc2;
    data1 ^= data3;  crc1 ^= crc3;
    data2 ^= data3;  crc2 ^= crc3;
    data3 ^= d;  crc3 ^= c;
    uint32_t poly = 1;
    do 
        if (calculate_crc(data1, poly, 0) == crc1 &&
            calculate_crc(data2, poly, 0) == crc2 &&
            calculate_crc(data3, poly, 0) == crc3)
            printf("poly = %08x, xor = %08x\n",
                   poly, calculate_crc(d, poly, 0) ^ c);
        poly += 2;
     while (poly != 1);


int main(void) 
    findp(0x928F640C, 0x0121B30E, 0xCB652607,
          0x70764315, 0x1F716284, 0x83692678);
    return 0;

通过求解 GF(2) 上的一组线性方程，有一种更快的方法，实际上速度更快。但是，我需要超过 40 秒的时间来编写该代码，所以这就是我要停下来的地方。除非我有很多很多这样的 CRC 可以找到。或者，除非我试图找到一个 64 位 CRC 多项式。

【讨论】：

效果很好！非常感谢您的提示马克！ :)

@MarkAdler - 只有一对样本异或在一起，问题是一个方程，两个变量（多项式和商），因此需要更多样本。