Суть в чем:
Есть очень коллизийный алго хэширования:
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] = {};
Полный код скрипта:
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;
}