生成 CUDA __device__ 内核内的 char 数组的所有组合

Question

我需要帮助。我开始使用 CUDA (2.3 / 3.0beta) 编写一个常见的暴力破解器/密码猜测器。 我尝试了不同的方法来生成定义的 ASCII 字符集的所有可能的纯文本“候选者”。

在此示例代码中，我想生成所有 74^4 种可能的组合（并将结果输出回主机/标准输出）。

$ ./combinations
Total number of combinations    : 29986576

Maximum output length   : 4
ASCII charset length    : 74

ASCII charset   : 0x30 - 0x7a
 "0123456789:;<=>?@ABCDEFGHIJKLMNOPQRSTUVWXYZ[\]^_`abcdefghijklmnopqrstuvwxy"

CUDA 代码（使用 2.3 和 3.0b 编译 - sm_10）-combinaions.cu：

#include <stdio.h>
#include <cuda.h>

__device__ uchar4 charset_global = {0x30, 0x30, 0x30, 0x30};
__shared__ __device__ uchar4 charset[128];

__global__ void combo_kernel(uchar4 * result_d, unsigned int N)
{
 int totalThreads = blockDim.x * gridDim.x ;
 int tasksPerThread = (N % totalThreads) == 0 ? N / totalThreads : N/totalThreads + 1;
 int myThreadIdx = blockIdx.x * blockDim.x + threadIdx.x ;
 int endIdx = myThreadIdx + totalThreads * tasksPerThread ;
 if( endIdx > N) endIdx = N;

 const unsigned int m = 74 + 0x30;

 for(int idx = myThreadIdx ; idx < endIdx ; idx += totalThreads) {
  charset[threadIdx.x].x = charset_global.x;
  charset[threadIdx.x].y = charset_global.y;
  charset[threadIdx.x].z = charset_global.z;
  charset[threadIdx.x].w = charset_global.w;
  __threadfence();

  if(charset[threadIdx.x].x < m) {
   charset[threadIdx.x].x++;

  } else if(charset[threadIdx.x].y < m) {
   charset[threadIdx.x].x = 0x30; // = 0
   charset[threadIdx.x].y++;

  } else if(charset[threadIdx.x].z < m) {
   charset[threadIdx.x].y = 0x30; // = 0
   charset[threadIdx.x].z++;

  } else if(charset[threadIdx.x].w < m) {
   charset[threadIdx.x].z = 0x30;
   charset[threadIdx.x].w++;; // = 0
  }

  charset_global.x = charset[threadIdx.x].x;
  charset_global.y = charset[threadIdx.x].y;
  charset_global.z = charset[threadIdx.x].z;
  charset_global.w = charset[threadIdx.x].w;

  result_d[idx].x = charset_global.x;
  result_d[idx].y = charset_global.y;
  result_d[idx].z = charset_global.z;
  result_d[idx].w = charset_global.w;
 }
}

#define BLOCKS 65535
#define THREADS 128

int main(int argc, char **argv)
{
 const int ascii_chars = 74;
 const int max_len = 4;
 const unsigned int N = pow((float)ascii_chars, max_len);
 size_t size = N * sizeof(uchar4);

 uchar4 *result_d, *result_h;
 result_h = (uchar4 *)malloc(size );
 cudaMalloc((void **)&result_d, size );
 cudaMemset(result_d, 0, size);

 printf("Total number of combinations\t: %d\n\n", N); 
 printf("Maximum output length\t: %d\n", max_len);
 printf("ASCII charset length\t: %d\n\n", ascii_chars);

 printf("ASCII charset\t: 0x30 - 0x%02x\n ", 0x30 + ascii_chars);
 for(int i=0; i < ascii_chars; i++)
  printf("%c",i + 0x30);
 printf("\n\n");

 combo_kernel <<< BLOCKS, THREADS >>> (result_d, N);
 cudaThreadSynchronize();

 printf("CUDA kernel done\n");
 printf("hit key to continue...\n");
 getchar();

 cudaMemcpy(result_h, result_d, size, cudaMemcpyDeviceToHost);

 for (unsigned int i=0; i<N; i++)
  printf("result[%06u]\t%c%c%c%c\n",i, result_h[i].x, result_h[i].y, result_h[i].z, result_h[i].w);

 free(result_h);
 cudaFree(result_d);
}

代码应该编译没有任何问题，但输出不是我所期望的。

在模拟模式上：

CUDA kernel done hit
key to continue...

    result[000000]  1000 
...
    result[000128]  5000

在发布模式上：

CUDA kernel done hit
key to continue...

    result[000000]  1000 
...
    result[012288]  5000

我还在代码的不同行上使用了 __threadfence() 和/或 __syncthreads() 也没有成功...

ps。如果可能的话，我想生成内核函数内的所有内容。我还尝试在 host main 函数和 memcpy 到 device 中“预”生成可能的纯文本候选，这只适用于非常有限的字符集大小（因为设备内存有限） ）。

• 关于输出的任何想法，为什么重复（即使使用 __threadfence() 或 __syncthreads()）？

  关于

• 任何其他方法可以在 CUDA 内核中快速生成纯文本（候选）:-) (~75^8)？

谢谢一百万

问候一月

Answer 1

顺便说一句，您的循环界限过于复杂。您不需要执行所有这些工作来计算 endIdx，而是可以执行以下操作，使代码更简单。

for(int idx = myThreadIdx ; idx < N ; idx += totalThreads)

Answer 2

让我们看看：

• 填充字符集数组时，__syncthreads() 就足够了，因为您对写入全局内存不感兴趣（稍后会详细介绍）
• 您的 if 语句不是正确重置循环迭代器：
  • 在 z < m，则 x == m 和 y == m 都必须设置为 0。
  • 与 w 类似
• 每个线程负责写入一组charset 中有 4 个字符，但每个线程都写入相同的 4 个值。没有线程做任何独立的工作。
• 您正在将每个线程结果写入全局内存，而没有原子操作，这是不安全的。无法保证结果在读回之前不会立即被另一个线程破坏。
• 将计算结果写入全局内存后，您将立即从全局内存中读回计算结果。目前还不清楚你为什么要这样做，而且这是非常不安全的。
• 最后，CUDA 中没有可靠的方法来实现所有块之间的同步，这似乎正是您所希望的。调用 __threadfence 仅适用于当前在设备上执行的块，该块可以是应为内核调用运行的所有块的子集。因此它不能作为同步原语工作。

计算每个线程的 x、y、z 和 w 的初始值可能更容易。然后每个线程可以从其初始值开始循环，直到执行完tasksPerThread 迭代。写出这些值可能或多或少可以像您现在所拥有的那样进行。

编辑：这是一个简单的测试程序，用于演示循环迭代中的逻辑错误：

int m = 2;
int x = 0, y = 0, z = 0, w = 0;

for (int i = 0; i < m * m * m * m; i++)
{
    printf("x: %d y: %d z: %d w: %d\n", x, y, z, w);
    if(x < m) {
        x++;
    } else if(y < m) {
        x = 0; // = 0
        y++;
    } else if(z < m) {
        y = 0; // = 0
        z++;
    } else if(w < m) {
        z = 0;
        w++;; // = 0
    }
}

其输出如下：

x: 0 y: 0 z: 0 w: 0
x: 1 y: 0 z: 0 w: 0
x: 2 y: 0 z: 0 w: 0
x: 0 y: 1 z: 0 w: 0
x: 1 y: 1 z: 0 w: 0
x: 2 y: 1 z: 0 w: 0
x: 0 y: 2 z: 0 w: 0
x: 1 y: 2 z: 0 w: 0
x: 2 y: 2 z: 0 w: 0
x: 2 y: 0 z: 1 w: 0
x: 0 y: 1 z: 1 w: 0
x: 1 y: 1 z: 1 w: 0
x: 2 y: 1 z: 1 w: 0
x: 0 y: 2 z: 1 w: 0
x: 1 y: 2 z: 1 w: 0
x: 2 y: 2 z: 1 w: 0