CUDA Source-code

2011-06-08T11:31:00.002+02:00

#include<stdio.h>

#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.

Prinzip des Algorithmus

2011-06-06T12:56:00.001+02:00

Das Prinzip des Algorithmus ist in allen Programmiersprachen gleich.

Ein Zähler zählt jede Runde um eins hinauf. Diese generierte Zahl wird dann in die einzelnen Buchstaben des zu generierenden Passworts umgewandelt, zusammengefügt ergeben sie eine Buchstabenkombination, welche das gesuchte Passwort repräsentieren. Es werden alle Buchstabenkombinationen nach der Reihe generiert, beginnend bei einer Passwortlänge von eins bis zu einer maximalen Passwortlänge von fünf, beziehungsweise von sieben.

Jede Zahl steht repräsentativ für einen Buchstaben, beziehungsweise für ein Passwort:

a	b	c	d	e	f	g	h	i	j	k	l	m
0	1	2	3	4	5	6	7	8	9	10	11	12
n	o	p	q	r	s	t	u	v	w	x	y	z
13	14	15	16	17	18	19	20	21	22	23	24	25

Ein Zähler, wird bei jedem neu zu generierendem Passwort inkrementiert:

Passwort	abc	abd	abe	abf	abg	abh	abi	abj	abk
Zähler++	28	29	30	31	32	33	34	35	36

Ein wichtiger Faktor um aus einer Zahl das richtige Passwort zu errechnen, ist die Passwortlänge. Das Passwort ist beim Starten eine Stelle lang, und generiert somit alle möglichen Kombinationen mit einer Stelle: a-z.
Jede Runde wird die Passwortlänge um eins erhöht.
In der zweiten Runde werden alle Kombinationen mit zwei Stellen generiert: aa-zz. In der dritten Runde: aaa-zzz. Bis zu maximal sieben Stellen, können alle Kombinationen errechnet werden.

Umrechnungsbeispiel:

· Gegeben ist die Zahl 3.206.568.

· Die Passwortlänge beträgt 5.

· Die Alphabetlänge beträgt 26 (nur Kleinbuchstaben)

Man nimmt die Zahl und dividiert sie durch die Aplphabetlänge.
Die Restklasse ist der 1. Ausgerechnete Buchstabe.
Das abgerundete Ergebnis dient zu ausrechnen des zweiten Buchstaben.

Dieser Vorgang wird so oft wiederholt, wie die Passwortlänge lang ist.

In diesem Beispiel ist das gesuchte Passwort 5 Buchstaben lang, somit wird der Vorgang 5 Mal wierderholt:

Runde	Rechnung	Ergebnis	Rest	Buchstabe
1	3.206.568 / 26	123.329	14	O
2	123.329 / 26	4.743	11	L
3	4.743 / 26	182	11	L
4	182 / 26	7	0	A
5	7 / 26		7	H

Der Rest jeder Rechnung steht für einen Buchstaben, zusammengefügt ergibt es das Passwort: HALLO

Vergleich zwischen CPU und GPU

2011-06-06T11:19:00.000+02:00

Benchmarktest C# Brute Force

2011-03-29T08:09:00.009+02:00

Die Grafik stellt die Relation von Zeitaufwand und Anzahl der Threads dar.

Beginnend mit 100 Threads, wurden jeder Runde 100 Threads hinzugefügt.

Als Suchwort dient "abcdefg".

Man sieht, dass sich die Geschwindigkeit steigert, jedoch nur bis zu einer gewissen Anzahl an Threads. Diese Anzahl der sich rentierenden Threads will ich anhand von Benchmarktests ermitteln.

Facebook-Fan/Infopage erstellt!

2011-02-22T09:38:00.002+01:00

Wir haben heute eine Facebook-Seite erstellt.
Ihr könnt auf dieser kommentieren oder uns eine Freundschaftsanfrage schicken.

Viel Spaß.

Herzlich Wilkommen!

2011-02-15T09:37:00.001+01:00

Hier werden im Laufe der Zeit, Informationen über einen Brute-Force-Angriff unter CUDA bereitgestellt.

a	b	c	d	e	f	g	h	i	j	k	l	m
0	1	2	3	4	5	6	7	8	9	10	11	12
n	o	p	q	r	s	t	u	v	w	x	y	z
13	14	15	16	17	18	19	20	21	22	23	24	25

a	b	c	d	e	f	g	h	i	j	k	l	m
0	1	2	3	4	5	6	7	8	9	10	11	12
n	o	p	q	r	s	t	u	v	w	x	y	z
13	14	15	16	17	18	19	20	21	22	23	24	25

Nvidia CUDA - HTL Ottakring

CUDA Source-code

Prinzip des Algorithmus

Vergleich zwischen CPU und GPU

Benchmarktest C# Brute Force

Facebook-Fan/Infopage erstellt!

Herzlich Wilkommen!

a	b	c	d	e	f	g	h	i	j	k	l	m
0	1	2	3	4	5	6	7	8	9	10	11	12
n	o	p	q	r	s	t	u	v	w	x	y	z
13	14	15	16	17	18	19	20	21	22	23	24	25