Помощь в оптимизации хэш-брута

[DrMefistO]

Доброго времени суток! Я новичок в CUDA-вычислениях, но то, что я успел изучить, помогло мне написать алгоритм перебора. Видюха используется NVidia GeForce GTX 450.
Суть в чем:
Есть очень коллизийный алго хэширования:

Code: Select all

__device__ unsigned int sdbm(const char * str, unsigned int init){
	unsigned int hash = init;
	int i = 0;	
	while (str[i]) hash = 0x1003F * (hash + str[i++]);
	return hash;
}

И есть два хэша:
- Хэш полного пути к файлу, без расширения (пример: render/stateblocks/default = 0xA9598550);
- Хэш имени файла с расширением (пример: default.rsb = 0x5134CB8C)

Так вот, у меня в скрипте определен массив результатов, где хранятся текущие и, они же, конечные результаты подсчетов:

Code: Select all

#define MAX_THREADS 960
#define MAX_BLOCKS 85
#define MAX_PATH 10
#define MAX_PATHS (MAX_BLOCKS * MAX_THREADS)

__device__ char paths[MAX_PATHS][MAX_PATH] = {};

Есть подозрение, что работа с данным массивом не такая быстрая, как могла быть. И, вторая проблема: повышение кол-ва тредов дает cudaErrorLaunchOutOfResources, а повышение кол-ва блоков дает StackOverflowError. В чем может быть беда, и как ее решить?

Полный код скрипта:

Code: Select all

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>
#include "string.h"
#include <cstdlib>

#define MAX_THREADS 960
#define MAX_BLOCKS 85
#define MAX_PATH 10
#define MAX_PATHS (MAX_BLOCKS * MAX_THREADS)

cudaError_t sdbmWithCuda();

__constant__ unsigned int prefix = 0x685D1304;
__constant__ unsigned int start = 0xEA298DB3;
__constant__ unsigned int hash1 = 0x620CB43B;
__constant__ unsigned int hash2 = 0xB9C2E077;
__constant__ char ext[] = ".rsb";
__constant__ char alpha[] = "abcdefghiklmnopqrstuvwxy_";
__constant__ int alpha_len = sizeof(alpha) - 1;
__device__ unsigned long long passed = 0x0;
__device__ int founded[1];
__device__ char paths[MAX_PATHS][MAX_PATH] = {};

__device__ unsigned int sdbm(const char * str, unsigned int init){
	unsigned int hash = init;
	int i = 0;	
	while (str[i]) hash = 0x1003F * (hash + str[i++]);
	return hash;
}

__device__ unsigned int sdbm_ext(const char * str){
	int i = 0;
	unsigned int init = sdbm(str, prefix);
	while (ext[i]) init = 0x1003F * (init + ext[i++]);
	return init;
}

__device__ int idx2len(unsigned long long idx){
	int len = 1;
	
	while (idx >= alpha_len){
		idx = (idx / alpha_len) - 1;
		len++;
	}
	
	return len;
}

__device__ int idx2name(unsigned long long idx, char * path){
	int len = idx2len(idx);
	int retn = len;
	path[len] = '\0';

	path[--len] = alpha[idx % alpha_len];
	while (idx >= alpha_len){
		idx = (idx / alpha_len) - 1;
		path[--len] = alpha[idx % alpha_len];
	}
	return retn;
}

__global__ void sdbmKernel(unsigned long long passed_)
{ 
	int thread = threadIdx.x + blockIdx.x * blockDim.x;
	passed = passed_;
	unsigned long long idx = thread + passed;	

	idx2name(idx, paths[thread]);

	unsigned int hash1_ = sdbm(paths[thread], start);
	unsigned int hash2_ = sdbm_ext(paths[thread]);

	if (hash1_ != hash1 || hash2_ != hash2){
		paths[thread][0] = '\0';
	}
	else if (hash1_ == hash1 && hash2_ == hash2)
		founded[0]++;
}

int main(int argc, char *argv[])
{	
    cudaStatus = sdbmWithCuda();
    cudaStatus = cudaDeviceReset();
    return 0;
}

cudaError_t sdbmWithCuda()
{
    cudaError_t cudaStatus;
    char paths_[MAX_PATHS][MAX_PATH];
    char p[MAX_PATH];

    int i, founded_ = 0;
    unsigned long long idx = 0;
    FILE * f;

    cudaStatus = cudaSetDevice(0);

    f = fopen("log_strings.txt", "wb+");

    while (founded_ < 1){
        sdbmKernel<<<MAX_BLOCKS, MAX_THREADS>>>(idx);
	idx += MAX_BLOCKS * MAX_THREADS;
    
        cudaDeviceSynchronize();
	cudaMemcpyFromSymbol(&founded_, founded, sizeof(int), 0, cudaMemcpyDeviceToHost);

	if (!founded_) continue;

	cudaStatus = cudaMemcpyFromSymbol(paths_, paths, sizeof(paths));

	for (i = 0; i < MAX_PATHS; i++)
	    if (strlen(paths_[i])){
	        printf("\n\n\r%s", paths_[i]);
		fwrite(paths_[i], strlen(paths_[i]), 1, f);
		fwrite("\n", 1, 1, f);
	    }
	}
	
	cudaFree(paths);
	fclose(f);

    return cudaStatus;
}

[DrMefistO]

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