1
我需要帮助。我开始用CUDA（2.3/3.0beta）编写一个常见的暴力破解/密码猜测器。 我尝试了不同的方法来生成一个定义的ASCII字符集的所有可能的纯文本“候选者”。在CUDA __device__内核中生成char数组的所有组合
在这个示例代码中，我想要生成所有74^4种可能的组合（并只将结果输出回主机/标准输出）。
$ ./combinations
Total number of combinations : 29986576
Maximum output length : 4
ASCII charset length : 74
ASCII charset : 0x30 - 0x7a
":;<=>[email protected][\]^_`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。如果可能的话，我想生成内核函数中的所有内容。我还尝试过在“主机主要功能和memcpy内部产生可能的纯文本候选者”，这只适用于非常有限的字符集大小（由于设备内存有限）。
有关输出的任何想法，为什么重复（甚至用__threadfence（）或__syncthreads（））？
任何其他在CUDA内核中快速生成纯文本（候选）的方法:-)（〜75^8）？
万分感谢
迎接一月
嗨，谢谢你的回答！ 我的想法是“__device__ uchar4 charset_global”是一种主数组。 每个线程块都应该将“charset_global的当前值”提取到共享字符集[128]中，做下一个组合（用这里的char设置填充一些计算），最后将“已由线程计算出来”组合写入charset_global var 。 （所以下一个线程可以使用“已完成组合”作为偏移量）。 我希望你对我有帮助;）） ps。 “你的if语句没有正确地重置你的循环迭代器” - 应该正确的工作在userland上 - origin：combfunc aocp – sead
我不知道什么是'在userland中正确工作'的意思，但你可以看到在我的编辑中使用代码循环迭代确实存在问题。 – Eric
您所描述的算法（在您的评论中）是一种串行算法。也就是说，没有线程可以计算唯一的密码，直到它从前一个线程获得结果。没有线程可以并行操作，因为它们将以相同的初始密码开始并以相同的方式进行排列，从而产生重复的输出。并行化的方法是理解您将生成74^N个可能的组合，每个线程将生成74^N/M个完全独立于其他线程的74^N/M个组合。 – Eric