#include<stdlib.h>
#include<string.h>
#include <time.h>
#include "cuda.h"
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#define EINGABELAENGE 8
//Da vom Kernelcode keine Hostmethoden aufgerufen werden können werden alle Methoden
//die im Kernel gebraucht werden kopiert und als __Device__ Methoden deklariert
//__device__ methoden können nur von anderen __device__ methoden oder vom Kernel //aufgerufen werden
__device__
char *cudaStrcpy(char *dest, const char *src)
{
unsigned i;
for (i=0; src[i] != '\0'; ++i)
dest[i] = src[i];
dest[i] = '\0';
return dest;
}
__device__
int cudaStrcmp(char *gesamtString, char *gesWort, int intervall)
{
int i;
for (i = 0; gesamtString[i] == gesWort[i]; i++)
if (i == intervall - 1 && gesWort[i + 1] == '\0')
return 0;
return 1;
}
__device__
long cudaPow(int z, int h)
{
int i;
long temp = 1;
for(i = 0; i < h; i++)
{
temp *=z;
}
return temp;
}
long pow(int z, int h)
{
int i;
long temp = 1;
for(i = 0; i < h; i++)
{
temp *=z;
}
return temp;
}
//Kernel = Devicecode der vom Host aufgerufen wird, wird als __global__ deklariert
__global__
void BruteForce(char* charset, int wordlength, int charsetlenght, long* outp,
char* suche, char* s, long* d, int interval,
int offset, int round)
{
//tid ist die ThreadID da die Threads in CUDA mehrdimensional angegeben werden müssen //sie zuerst liniear gemacht werden
int tid = blockIdx.x * blockDim.x + threadIdx.x;
//Startwert und Endwert wird berechnet mit ThreadID * Interval
//für spätere Durchgänge, die mehr als 5 sekunden Rechenzeit bräuchten, muss die
//Aufgabe in kleinere Teile geteilt werden was mithilfe round und offset geschieht
long startw = round * offset + tid * interval;
long endw = round * offset + (tid + 1) * interval;
int i = 0;
long mw = 0;
long w = 0;
while(startw < endw)
{
//Die logik die aus Zahlen Buchstabenketten macht
mw = startw;
for(i = wordlength; i >= 0; i--)
{
w = (int)(mw / d[i]);
if ( i == wordlength)
{
if (w != 0)
{
s[tid * wordlength + i]=charset[w];
}
}
else
{
s[tid * wordlength + i] = charset[w];
}
mw = mw - w * d[i];
}
//Überprüft ob die aktuelle Buchstabenfolge die gesuchte ist
if(cudaStrcmp(&s[tid * wordlength],suche, wordlength) == 0)
{
outp[tid] = startw;
}
startw++;
}
}
// gibt den Startzeitpunkt und Endzeitpunkt aus
int printTime(tm startTime, tm endTime)
{
printf("\n\n%d:%d\n",startTime.tm_min,startTime.tm_sec);
printf("%d:%d\n",endTime.tm_min,endTime.tm_sec);
return 0;
}
//Überprüft ob ein Thread ein ergebnis zurückliefert
int checkResults(long* results, int resultCount)
{
int c1 = 0;
for(c1 = 0; c1 < resultCount; c1++)
{
if(results[c1] != -1)
{
printf("\n%d:%d",c1,results[c1]);
return results[c1];
}
}
return -1;
}
int main(int argc, char** argv)
{
time_t rawtime;
struct tm * timeinfo;
struct tm tmstart;
struct tm tmend;
int gefunden = -1;
int teiler = 0;
int offset = 0;
//h_ vor dem Variablenname bedeutet das es sich um eine hostvariable handelt
long * h_outp;
int h_i = 0;
int h_alphabetlaenge = 26;
char h_loweralpha[27] = "abcdefghijklmnopqrstuvwxyz";
char* h_eingabe;
long* h_d;
int h_interval;
int h_threadanzahl;
int c1 = 0;
h_eingabe = (char*)malloc(sizeof(char)*EINGABELAENGE);
h_outp = (long*)malloc(sizeof(long) * 26);
//d_ vor dem Variablenname bedeutet das es sich um eine devicevariable handelt
long* d_outp;
char* d_loweralpha;
char* d_eingabe;
char* d_s;
long* d_d;
//Speicherzuweisung von devicespeicher für die Variablen die bei jeden durchgang
//gleich bleiben
cudaMalloc((void **) &d_eingabe,sizeof(char)*EINGABELAENGE);
cudaMalloc((void **) &d_loweralpha,sizeof(char)*h_alphabetlaenge);
cudaMalloc((void **) &d_outp,sizeof(long) * 26);
//Das zu suchende Wort wird eingegeben
scanf("%s",h_eingabe);
//Der Startzeitpunkt wird festgelegt und in tmstart gespeichert
time ( &rawtime );
timeinfo = localtime ( &rawtime );
tmstart = *timeinfo;
//Die Eingabe und das Alphabet wird auf devicespeicher kopiert
cudaMemcpy(d_eingabe,h_eingabe,EINGABELAENGE,cudaMemcpyHostToDevice);
cudaMemcpy(d_loweralpha,h_loweralpha,h_alphabetlaenge,cudaMemcpyHostToDevice);
//Die schleife läuft solange bis ein ergebnis gefunden wurde oder bis EINGABELAENGE
//erreicht ist wobei EINGABELAENGE sich auf die maximale und nicht auf die laenge der
// aktuellen eingabe bezieht (siehe macro definition)
for (h_i = 1; h_i < EINGABELAENGE && gefunden == -1; h_i++)
{
printf("\n\n");
//Es wird der Interval und die Threadanzahl bestimmt da es beim ersten durchgang
//nur soviele Durchläufe gibt wie das alphabet lang ist werden nicht mehr
//threads gebraucht
if(h_i == 1)
{
h_interval = 1;
h_threadanzahl = h_alphabetlaenge;
}
else
{
//die optimale Anzahl der Threads ist für jede Grafikkarte anders
//und lässt sich über CUDA devicequerry herausfinden
h_threadanzahl = 4 * 112;
h_interval = (pow(h_alphabetlaenge,h_i)/h_threadanzahl);
}
//Jeder thread schreibt in output rein ob und an welcher stelle er den gesuchten
// String gefunden hatt
h_outp = (long*)malloc(sizeof(long) * h_threadanzahl);
cudaMalloc((void **) &d_outp,sizeof(long) * h_threadanzahl);
//outp wird auf -1 initialisiert weil die werte die zurückgegeben werden
//theoretisch zwischen 0 und unendlich liegen können
for(c1 = 0; c1 < h_threadanzahl; c1++)
h_outp[c1] = -1;
//die Werte von h_outp(-1) werden auf devicememory kopiert
cudaMemcpy(d_outp,h_outp,sizeof(long) * h_threadanzahl,cudaMemcpyHostToDevice);
//in d wird der Teiler berechnet die zur berechnung der strings im
//Bruteforce algorithmus gebraucht werden
cudaMalloc((void **) &d_d,sizeof(long) * h_i);
//in “s“ werden die strings gespeichert die gerade überprüft werden
//da diese mit jedem durchlauf(i) um eine stelle größer werden und
//jeder thread seinen eigenen speicher braucht muss entsprechend
//viel speicher allokiert werden
cudaMalloc((void **) &d_s,sizeof(char*) * h_threadanzahl * h_i);
//die teiler werden berechnet
h_d = (long*)malloc(sizeof(long) * h_i);
for(c1 = h_i; c1 >= 0; c1--)
{
h_d[c1]= pow(h_alphabetlaenge,c1);
}
//und auf das device kopiert
cudaMemcpy(d_d, h_d, h_i * sizeof(int),cudaMemcpyHostToDevice);
//Diese if/else anweisung ist ein workaround um den „5 sekunden bug“ zu beheben
//wenn eine Kernelfunktion länger als ca. 5 sekunden zum abarbeiten braucht
//wird diese vom Windows Treiber unterbrochen
if(h_i > 4)
{
//das passwort länger als 4 zeichen ist muss die aufgabe geteilt werden
//dadurch das der aufwand exponentiell steigt steigt auch die anzahl
//der benötigten teilungen exponentiell
//diese aufteilung führt dazu das der kernel immer gleichlang ausgeführt wird
teiler = 2 * pow(h_alphabetlaenge,h_i-5);
h_interval = (pow(h_alphabetlaenge,h_i)/h_threadanzahl)/teiler + 1;
offset = pow(h_alphabetlaenge ,h_i)/teiler;
for(c1 = 0; c1 < teiler; c1++)
{
//kernelaufruf 4 ist die anzahl der blocks und treadanzahl/4 die anzahl
//der Threads pro block
BruteForce<<<4,h_threadanzahl/4>>>(d_loweralpha, h_i, h_alphabetlaenge,
d_outp, d_eingabe, d_s, d_d, h_interval, offset, c1);
printf(".");
//zwischenergebnis wird aus dem device kopiert
cudaMemcpy(h_outp,d_outp,sizeof(long)*
h_threadanzahl,cudaMemcpyDeviceToHost);
//und überprüft
gefunden = checkResults(h_outp,h_threadanzahl);
//falls ein ergebnis vorliegt wird die endzeit angezeigt
if(gefunden != -1)
{
time ( &rawtime );
timeinfo = localtime ( &rawtime );
tmend = *timeinfo;
printTime(tmstart,tmend);
break;
}
}
}
else
{
//wenn die wortlänge kleiner als 5 ist reicht ein kernellaufruf
//zum berechnen
offset = 0;
c1 = 0;
BruteForce<<<4,h_threadanzahl/4>>>(d_loweralpha, h_i, h_alphabetlaenge,
d_outp, d_eingabe, d_s, d_d, h_interval, offset, c1);
cudaMemcpy(h_outp,d_outp,sizeof(long) *
h_threadanzahl,cudaMemcpyDeviceToHost);
gefunden = checkResults(h_outp,h_threadanzahl);
if(gefunden != -1)
{
time ( &rawtime );
timeinfo = localtime ( &rawtime );
tmend = *timeinfo;
printTime(tmstart,tmend);
break;
}
// . geben die anzahl der ausgeführten kernelaufrufe an
printf(".");
}
}
free(h_outp);
free(h_eingabe);
cudaFree(d_d);
cudaFree(d_s);
cudaFree(d_outp);
cudaFree(d_loweralpha);
cudaFree(d_eingabe);
cudaFree(d_s);
system("PAUSE");
return 0;
}
Jeder “.“ steht für einen abgeschlossenen Kernelaufruf, jede neue Zeile für eine neue Wortlänge die überprüft wird. 86 gibt die ThreadID des Threads an der den Treffer gelandet hat und 268478924 sind die Anzahl der versuche die es gebraucht hatt das Ergebnis zu finden.
