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