#include <stdlib.h>
#include <stdio.h>

#define output 1

//Variablen
#define p 5
#define q 7
#define n 35
#define e 5
#define v 5
#define z 24
#define anzahl_Zeichen 2688
#define count_cores 384

/*
Klartext: K
Geheimtext: G
Verschluesselung: G = K^v mod n
Entschluesselung: K = G^e mod n
Index des CUDA Kerns: blockIdx.x blockIdx.y
Ein groesserer Text soll ver- und entschluesselt werden.
Dieser wird jedoch wie folgt veraendert: nur kleine Buchstaben, keine Sonderzeichen außer . und ,
Dafuer werden die einzellnen chars in Integer umgewandelt (<=30).
Somit ist eine Verarbeitung moeglich.
*/

static void HandleError( cudaError_t err, const char *file, int line ) {
	if (err != cudaSuccess) {
		printf( "%s in %s at line %d\n", cudaGetErrorString( err ), file, line );
		exit( EXIT_FAILURE );
	}
}

#define HANDLE_ERROR( err ) (HandleError( err, __FILE__, __LINE__ ))

__global__ void verschluessselung(long int klartexte[], long int geheimtexte[])
{
	long int j, multi, x;
	
	//Fuer 384 Bloecke mit Threads
	long int threads = anzahl_Zeichen/count_cores;

	multi = x  = klartexte[threadIdx.x+blockIdx.x*threads];
	for (j = 1; j < v; j++)
		x *= multi;
	
	geheimtexte[threadIdx.x+blockIdx.x*threads] = x % n;
}

__global__ void entschluessselung(long int geheimtexte[], long int klartexte_pruefung[])
{
	long int j, multi, x;
	
	//Fuer 384 Bloecke mit Threads
	long int threads = anzahl_Zeichen/count_cores;
	
	multi = x  = geheimtexte[threadIdx.x+blockIdx.x*threads];
	for (j = 1; j < e; j++)
		x *= multi;
	
	klartexte_pruefung[threadIdx.x+blockIdx.x*threads] = x % n;
}

/*
 * Umwandlung von char in Zahl 0 bis 30 .
 */
void splitt(char text[], long int numbers[])
{
	int i;
	
	//Splitte Klartext
	for (i = 0; i < anzahl_Zeichen; i++)
	{
		long int number = (int)text[i];
		
		//char in int beginnend mit 0
		//Sonderzeichen
		if (number == 44)		//,
			number = 27;
		else if (number == 46)		//.
			number = 28;
		else if (number == 10)		//\n
			number = 29;
		else if (number == 32)		//' '
			number = 30;
		else				//a-z
			number -= 97;
		
		numbers[i] = number;
	}
}

/*
 * Umwandlung von Zahl 0 bis 30 zu char.
 */
void unsplitt(char text[], long int numbers[])
{
	int i;
	
	//Splitte Klartext
	for (i = 0; i < anzahl_Zeichen; i++)
	{
		long int number = numbers[i];
		char t;
		
		//int in char
		//Sonderzeichen
		if (number == 27)		//,
			t = ',';
		else if (number == 28)		//.
			number = t = '.';
		else if (number == 29)		//\n
			t = '\n';
		else if (number == 30)		//' '
			t = ' ';
		else				//a-z
			t = (char)(number+97);
			
		text[i] = t;
	}
}

int main(int argc, char *argv[]) {
	int i;
	cudaEvent_t start, stop;
	float elapsedTime;
	int count_Threads;
	long int multi = 0;
	
	char klartext[anzahl_Zeichen+1];
	char klartext2[anzahl_Zeichen+1];
	long int kt_splitted[anzahl_Zeichen+1];
	long int kt_splitted2[anzahl_Zeichen+1];
	long int *dev_kt_splitted, *dev_kt_splitted2, *dev_gt_splitted;
	int size = sizeof(long int)*(anzahl_Zeichen+1);
	
	printf("\n-|| RSA mit CUDA ||-\n\n\n");
	
	//Klartetext erzeugen
	klartext[anzahl_Zeichen] = klartext2[anzahl_Zeichen] = '\0';
	strcpy(klartext, "hat der alte hexenmeister?sich doch einmal wegbegeben.?und nun sollen seine geister?auch nach meinem willen leben.?seine wort und werke?merkt ich und den brauch,?und mit geistesstaerke?tu ich wunder auch.?walle. walle?manche strecke,?dass, zum zwecke,?wasser fliesse?und mit reichem, vollem schwalle?zu dem bade sich ergiesse.?und nun komm, du alter besen.?nimm die schlechten lumpenhuellen ?bist schon lange knecht gewesen:?nun erfuelle meinen willen.?auf zwei beinen stehe,?oben sei ein kopf,?eile nun und gehe?mit dem wassertopf.?walle. walle?manche strecke,?dass, zum zwecke,?wasser fliesse?und mit reichem, vollem schwalle?zu dem bade sich ergiesse.?seht, er laeuft zum ufer nieder,?wahrlich. ist schon an dem flusse,?und mit blitzesschnelle wieder?ist er hier mit raschem gusse.?schon zum zweiten male.?wie das becken schwillt.?wie sich jede schale?voll mit wasser fuellt.?stehe. stehe.?denn wir haben?deiner gaben?vollgemessen. ?ach, ich merk es. wehe. wehe.?hab ich doch das wort vergessen.?ach, das wort, worauf am ende?er das wird, was er gewesen.?ach, er laeuft und bringt behende.?waerst du doch der alte besen.?immer neue guesse?bringt er schnell herein,?ach. und hundert fluesse?stuerzen auf mich ein.?nein, nicht laenger?kann ichs lassen ?will ihn fassen.?das ist tuecke.?ach. nun wird mir immer baenger.?welche mine. welche blicke.?o du ausgeburt der hoelle.?soll das ganze haus ersaufen??seh ich ueber jede schwelle?doch schon wasserstroeme laufen.?ein verruchter besen,?der nicht hoeren will.?stock, der du gewesen,?steh doch wieder still.?willst am ende?gar nicht lassen??will dich fassen,?will dich halten?und das alte holz behende?mit dem scharfen beile spalten.?seht da kommt er schleppend wieder.?wie ich mich nur auf dich werfe,?gleich, o kobold, liegst du nieder ?krachend trifft die glatte schaerfe.?wahrlich, brav getroffen.?seht, er ist entzwei.?und nun kann ich hoffen,?und ich atme frei.?wehe. wehe.?beide teile?stehn in eile?schon als knechte?voellig fertig in die hoehe.?helft mir, ach. ihr hohen maechte.?und sie laufen. nass und naesser?wirds im saal und auf den stufen.?welch entsetzliches gewaesser.?herr und meister. hoer mich rufen.  ?ach, da kommt der meister.?herr, die not ist gross.?die ich rief, die geister?werd ich nun nicht los.?in die ecke,?besen, besen.?seids gewesen.?denn als geister?ruft euch nur zu diesem zwecke,?erst hervor der alte meister.?                                                                                                                                                                                                                                                                                                    ");
	printf("\n\nDer Klartext ist %d Zeichen lang.\n", sizeof(klartext)/sizeof(char)-1);
	
	if (argc < 2)
	{
		printf("\nParameter fuer Groesse der Klartexted fehlt!\n");
		exit(0);
	}

	//lese Anzahl der Zeichen
	multi = atoi(argv[1]);

	if (multi < 1)
	{
		printf("\nAnzahl der Zeichen ist eine ganze positive Zahl.\n");
		exit(0);
	}
	
	printf("\nmulti: %d\n", multi);
	
	//Ausgabe
	if (output)
	{
		printf("\n\nAnfang des  Klartextes:\n\n");
		for (i = 0; i < anzahl_Zeichen; i++)
		{
			if (klartext[i] == '?')
				klartext[i] = '\n';
			if (i < 546)
				putchar(klartext[i]);
		}
		printf("\n\n");
	}
	
	//klartext2 mit a füllen
	for (i = 0; i < anzahl_Zeichen; i++)
	{
		klartext2[i] = 'a';
	}
	
	printf("Der Klartext wird nun verschluesselt und anschliessend entschluesselt.\n");
	
	//Chars in ints aufsplitten
	splitt(klartext, kt_splitted);
	
	//Variablen der Zeitmessung erstellen
	HANDLE_ERROR(cudaEventCreate(&start));
	HANDLE_ERROR(cudaEventCreate(&stop));
	
	//Start Zeitmessung
	HANDLE_ERROR(cudaEventRecord(start, 0));
	
	//allokieren
	HANDLE_ERROR(cudaMalloc((void **)&dev_kt_splitted, size));
	HANDLE_ERROR(cudaMalloc((void **)&dev_kt_splitted2, size));
	HANDLE_ERROR(cudaMalloc((void **)&dev_gt_splitted, size));
	
	//kopieren
	HANDLE_ERROR(cudaMemcpy(dev_kt_splitted, kt_splitted, size, cudaMemcpyHostToDevice));
	
	//Anzahl Threads pro Block
	count_Threads = anzahl_Zeichen/count_cores;
	
	for (i = 0; i < multi; i++)
	{
		//verschluesseln
		verschluessselung<<<count_cores, count_Threads>>>(dev_kt_splitted, dev_gt_splitted);
		
		//sync
		HANDLE_ERROR(cudaDeviceSynchronize());
		
		//entschluesseln
		entschluessselung<<<count_cores, count_Threads>>>(dev_gt_splitted, dev_kt_splitted2);
		
		//sync
		HANDLE_ERROR(cudaDeviceSynchronize());
	}
	
	//zurueckkopieren
	HANDLE_ERROR(cudaMemcpy(kt_splitted2, dev_kt_splitted2, size, cudaMemcpyDeviceToHost));
		
	//Ende der Zeitmessung
	HANDLE_ERROR(cudaEventRecord(stop, 0));
	HANDLE_ERROR(cudaEventSynchronize(stop));
	
	//Ausgabe der verstrichenen Zeit
	HANDLE_ERROR(cudaEventElapsedTime(&elapsedTime, start, stop));
	printf("\nBeendet.\n\n\nverstrichene Zeit: %3.1f ms\n", elapsedTime);
	
	//ints wieder in char umwandeln
	unsplitt(klartext2, kt_splitted2);
	
	//Ausgabe
	if (output)
	{
		printf("\n\nDer Klartext lautet nun: (Anfang)\n\n");
		for (i = 0; i < 545; i++)
		{
			putchar(klartext2[i]);
		}
		printf("\n\n");
	}
	
	//freigeben
	HANDLE_ERROR(cudaEventDestroy(start));
	HANDLE_ERROR(cudaEventDestroy(stop));
	HANDLE_ERROR(cudaFree(dev_kt_splitted));
	HANDLE_ERROR(cudaFree(dev_kt_splitted2));
	HANDLE_ERROR(cudaFree(dev_gt_splitted));
	return EXIT_SUCCESS;
}

	.file	"tmpxft_0026e2b9_00000000-6_rsa.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2033:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2033:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"/home/ubuntu/Datasets/stackv2/train-structured/TBoonX/krypto_cuda/master/rsa.cu"
.LC1:
	.string	"%s in %s at line %d\n"
	.text
	.type	_ZL11HandleError9cudaErrorPKci.constprop.0, @function
_ZL11HandleError9cudaErrorPKci.constprop.0:
.LFB2177:
	.cfi_startproc
	testl	%edi, %edi
	je	.L2
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	%esi, 12(%rsp)
	call	cudaGetErrorString@PLT
	movl	12(%rsp), %r8d
	movl	$2, %edi
	leaq	.LC0(%rip), %rcx
	movq	%rax, %rdx
	leaq	.LC1(%rip), %rsi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	$1, %edi
	call	exit@PLT
.L2:
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2177:
	.size	_ZL11HandleError9cudaErrorPKci.constprop.0, .-_ZL11HandleError9cudaErrorPKci.constprop.0
	.globl	_Z6splittPcPl
	.type	_Z6splittPcPl, @function
_Z6splittPcPl:
.LFB2028:
	.cfi_startproc
	endbr64
	xorl	%edx, %edx
.L12:
	movsbq	(%rdi,%rdx), %rax
	movl	$27, %ecx
	cmpb	$44, %al
	je	.L9
	jg	.L10
	movl	$29, %ecx
	cmpb	$10, %al
	je	.L9
	movl	$30, %ecx
	cmpb	$32, %al
	jmp	.L22
.L10:
	movl	$28, %ecx
	cmpb	$46, %al
.L22:
	je	.L9
	leaq	-97(%rax), %rcx
.L9:
	movq	%rcx, (%rsi,%rdx,8)
	incq	%rdx
	cmpq	$2688, %rdx
	jne	.L12
	ret
	.cfi_endproc
.LFE2028:
	.size	_Z6splittPcPl, .-_Z6splittPcPl
	.globl	_Z8unsplittPcPl
	.type	_Z8unsplittPcPl, @function
_Z8unsplittPcPl:
.LFB2029:
	.cfi_startproc
	endbr64
	xorl	%eax, %eax
	leaq	CSWTCH.44(%rip), %r8
.L26:
	movq	(%rsi,%rax,8), %rdx
	leaq	-27(%rdx), %r9
	leal	97(%rdx), %ecx
	cmpq	$3, %r9
	ja	.L25
	movb	-27(%r8,%rdx), %cl
.L25:
	movb	%cl, (%rdi,%rax)
	incq	%rax
	cmpq	$2688, %rax
	jne	.L26
	ret
	.cfi_endproc
.LFE2029:
	.size	_Z8unsplittPcPl, .-_Z8unsplittPcPl
	.globl	_Z39__device_stub__Z17verschluessselungPlS_PlS_
	.type	_Z39__device_stub__Z17verschluessselungPlS_PlS_, @function
_Z39__device_stub__Z17verschluessselungPlS_PlS_:
.LFB2055:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movq	%rdi, 8(%rsp)
	leaq	32(%rsp), %rcx
	leaq	24(%rsp), %rdx
	movq	%rsi, (%rsp)
	leaq	40(%rsp), %rdi
	leaq	52(%rsp), %rsi
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rax
	movl	$1, 48(%rsp)
	movq	%rax, 88(%rsp)
	movq	%rsp, %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 40(%rsp)
	movq	%rax, 52(%rsp)
	movl	$1, 60(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L28
	pushq	32(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z17verschluessselungPlS_(%rip), %rdi
	pushq	32(%rsp)
	.cfi_def_cfa_offset 144
	movq	68(%rsp), %rcx
	movl	76(%rsp), %r8d
	movq	56(%rsp), %rsi
	movl	64(%rsp), %edx
	leaq	104(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L28:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L30
	call	__stack_chk_fail@PLT
.L30:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2055:
	.size	_Z39__device_stub__Z17verschluessselungPlS_PlS_, .-_Z39__device_stub__Z17verschluessselungPlS_PlS_
	.globl	_Z17verschluessselungPlS_
	.type	_Z17verschluessselungPlS_, @function
_Z17verschluessselungPlS_:
.LFB2056:
	.cfi_startproc
	endbr64
	jmp	_Z39__device_stub__Z17verschluessselungPlS_PlS_
	.cfi_endproc
.LFE2056:
	.size	_Z17verschluessselungPlS_, .-_Z17verschluessselungPlS_
	.globl	_Z39__device_stub__Z17entschluessselungPlS_PlS_
	.type	_Z39__device_stub__Z17entschluessselungPlS_PlS_, @function
_Z39__device_stub__Z17entschluessselungPlS_PlS_:
.LFB2057:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movq	%rdi, 8(%rsp)
	leaq	32(%rsp), %rcx
	leaq	24(%rsp), %rdx
	movq	%rsi, (%rsp)
	leaq	40(%rsp), %rdi
	leaq	52(%rsp), %rsi
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rax
	movl	$1, 48(%rsp)
	movq	%rax, 88(%rsp)
	movq	%rsp, %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 40(%rsp)
	movq	%rax, 52(%rsp)
	movl	$1, 60(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L33
	pushq	32(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z17entschluessselungPlS_(%rip), %rdi
	pushq	32(%rsp)
	.cfi_def_cfa_offset 144
	movq	68(%rsp), %rcx
	movl	76(%rsp), %r8d
	movq	56(%rsp), %rsi
	movl	64(%rsp), %edx
	leaq	104(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L33:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L35
	call	__stack_chk_fail@PLT
.L35:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2057:
	.size	_Z39__device_stub__Z17entschluessselungPlS_PlS_, .-_Z39__device_stub__Z17entschluessselungPlS_PlS_
	.globl	_Z17entschluessselungPlS_
	.type	_Z17entschluessselungPlS_, @function
_Z17entschluessselungPlS_:
.LFB2058:
	.cfi_startproc
	endbr64
	jmp	_Z39__device_stub__Z17entschluessselungPlS_PlS_
	.cfi_endproc
.LFE2058:
	.size	_Z17entschluessselungPlS_, .-_Z17entschluessselungPlS_
	.section	.rodata.str1.1
.LC2:
	.string	"\n-|| RSA mit CUDA ||-\n\n\n"
.LC3:
	.ascii	"hat der alte hexenmeister?sich doch einmal wegbegeben.?und n"
	.ascii	"un sollen seine geister?auch nach meinem willen leben.?seine"
	.ascii	" wort und werke?merkt ich und den brauch,?und mit geistessta"
	.ascii	"erke?tu ich wunder auch.?walle. walle?manche strecke,?dass, "
	.ascii	"zum zwecke,?wasser fliesse?und mit reichem, vollem schwalle?"
	.ascii	"zu dem bade sich ergiesse.?und nun komm, du alter besen.?nim"
	.ascii	"m die schlechten lumpenhuellen ?bist schon lange knecht gewe"
	.ascii	"sen:?nun erfuelle meinen willen.?auf zwei beinen stehe,?oben"
	.ascii	" sei ein kopf,?eile nun und gehe?mit dem wassertopf.?walle. "
	.ascii	"walle?manche strecke,?dass, zum zwecke,?wasser fliesse?und m"
	.ascii	"it reichem, vollem schwalle?zu dem bade sich ergiesse.?seht,"
	.ascii	" er laeuft zum ufer nieder,?wahrlich. ist schon an dem fluss"
	.ascii	"e,?und mit blitzesschnelle wieder?ist er hier mit raschem gu"
	.ascii	"sse.?schon zum zweiten male.?wie das becken schwillt.?wie si"
	.ascii	"ch jede schale?voll mit wasser fuellt.?stehe. stehe.?denn wi"
	.ascii	"r haben?deiner gaben?vollgemessen. ?ach, ich merk es. wehe. "
	.ascii	"wehe.?hab ich doch das wort vergessen.?ach, das wort, worauf"
	.ascii	" am ende?er das wird, was er gewesen.?ach, er laeuft und bri"
	.ascii	"ngt behende.?waerst du doch der alte besen.?immer neue guess"
	.ascii	"e?bringt er schnell herein,?ach. und hundert fluesse?stuerze"
	.ascii	"n auf mich ein.?nein, nicht laenger?kann ichs lassen ?will i"
	.ascii	"hn fassen.?das ist tuecke.?ach. nun wird mir immer baenger.?"
	.ascii	"welche mine. welche blicke.?o du ausgeburt der hoelle.?soll "
	.ascii	"das ganze haus ersaufen??seh ich ueber jede schwelle?doch sc"
	.ascii	"hon wasserstroeme laufen.?ein verruchter besen,?der nicht ho"
	.ascii	"eren will.?stock, der du gewesen,?steh doch wieder still.?wi"
	.ascii	"llst am ende?gar nicht lassen??will dich fassen,?will dich h"
	.ascii	"alten?und das alte holz behende?mit dem scharfen beile spalt"
	.ascii	"en.?seht da kommt er schleppend wieder.?wie ich mich nur auf"
	.ascii	" dich werfe,?gleich, o kobold, liegst du nieder ?krachend tr"
	.ascii	"ifft die glatte schaerfe.?wahrlich, brav getroffen.?seht, er"
	.ascii	" ist entzwei.?und nun kann ich hoffen,?und ich atme frei.?we"
	.ascii	"he. wehe.?beide teile?stehn in eile?schon als knechte?voelli"
	.ascii	"g fertig in die hoeh"
	.ascii	"e.?helft mir, ach. ihr hohen maechte.?und sie laufen. nass u"
	.ascii	"nd naesser?wirds im saal und auf den stufen.?welch entsetzli"
	.ascii	"ches gewaesser.?herr und meister. hoer mich rufen.  ?ach, da"
	.ascii	" kommt der meister.?herr, die not ist gross.?die ich rief, d"
	.ascii	"ie geister?werd ich nun nicht los.?in die ecke,?besen, besen"
	.ascii	".?seids gewesen.?denn als geister?ruft euch nur zu diesem zw"
	.ascii	"ecke,?erst hervor der alte meister.?                        "
	.ascii	"            "
	.string	"                                                                                                                                                                                                                                                                "
.LC4:
	.string	"\n\nDer Klartext ist %d Zeichen lang.\n"
.LC5:
	.string	"\nParameter fuer Groesse der Klartexted fehlt!\n"
.LC6:
	.string	"\nAnzahl der Zeichen ist eine ganze positive Zahl.\n"
.LC7:
	.string	"\nmulti: %d\n"
.LC8:
	.string	"\n\nAnfang des  Klartextes:\n\n"
.LC9:
	.string	"\n\n"
.LC10:
	.string	"Der Klartext wird nun verschluesselt und anschliessend entschluesselt.\n"
.LC11:
	.string	"\nBeendet.\n\n\nverstrichene Zeit: %3.1f ms\n"
.LC12:
	.string	"\n\nDer Klartext lautet nun: (Anfang)\n\n"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2030:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	leaq	-45056(%rsp), %r11
	.cfi_def_cfa 11, 45112
.LPSRL0:
	subq	$4096, %rsp
	orq	$0, (%rsp)
	cmpq	%r11, %rsp
	jne	.LPSRL0
	.cfi_def_cfa_register 7
	subq	$3432, %rsp
	.cfi_def_cfa_offset 48544
	movq	%fs:40, %rax
	movq	%rax, 48472(%rsp)
	xorl	%eax, %eax
	movl	%edi, %ebp
	movq	%rsi, %rbx
	movl	$2, %edi
	leaq	.LC2(%rip), %rsi
	call	__printf_chk@PLT
	xorl	%eax, %eax
	movl	$2689, %ecx
	leaq	.LC3(%rip), %rsi
	leaq	43094(%rsp), %rdi
	movl	$2688, %edx
	movb	$0, 48471(%rsp)
	rep movsb
	leaq	.LC4(%rip), %rsi
	movl	$2, %edi
	call	__printf_chk@PLT
	decl	%ebp
	leaq	.LC5(%rip), %rsi
	jle	.L53
	movq	8(%rbx), %rdi
	call	atoi@PLT
	movslq	%eax, %rdx
	movq	%rdx, %r12
	testq	%rdx, %rdx
	jg	.L40
	leaq	.LC6(%rip), %rsi
.L53:
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	xorl	%edi, %edi
	call	exit@PLT
.L40:
	leaq	.LC7(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	xorl	%ebp, %ebp
	call	__printf_chk@PLT
	leaq	43094(%rsp), %r13
	leaq	.LC8(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	movq	%r13, %rbx
	call	__printf_chk@PLT
.L43:
	cmpb	$63, (%rbx)
	jne	.L41
	movb	$10, (%rbx)
.L41:
	cmpl	$545, %ebp
	jg	.L42
	movsbl	(%rbx), %edi
	call	putchar@PLT
.L42:
	incl	%ebp
	incq	%rbx
	cmpl	$2688, %ebp
	jne	.L43
	movl	$2, %edi
	xorl	%eax, %eax
	leaq	64(%rsp), %rbx
	xorl	%r15d, %r15d
	leaq	.LC9(%rip), %r14
	leaq	45783(%rsp), %rbp
	movq	%r14, %rsi
	call	__printf_chk@PLT
	movb	$97, %al
	movl	$2688, %ecx
	movq	%rbp, %rdi
	rep stosb
	leaq	.LC10(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	%rbx, %rsi
	movq	%r13, %rdi
	movabsq	$4294967680, %r13
	call	_Z6splittPcPl
	movq	%rsp, %rdi
	call	cudaEventCreate@PLT
	movl	$188, %esi
	movl	%eax, %edi
	call	_ZL11HandleError9cudaErrorPKci.constprop.0
	leaq	8(%rsp), %rdi
	call	cudaEventCreate@PLT
	movl	$189, %esi
	movl	%eax, %edi
	call	_ZL11HandleError9cudaErrorPKci.constprop.0
	movq	(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movl	$192, %esi
	movl	%eax, %edi
	call	_ZL11HandleError9cudaErrorPKci.constprop.0
	leaq	16(%rsp), %rdi
	movl	$21512, %esi
	call	cudaMalloc@PLT
	movl	$195, %esi
	movl	%eax, %edi
	call	_ZL11HandleError9cudaErrorPKci.constprop.0
	leaq	24(%rsp), %rdi
	movl	$21512, %esi
	call	cudaMalloc@PLT
	movl	$196, %esi
	movl	%eax, %edi
	call	_ZL11HandleError9cudaErrorPKci.constprop.0
	leaq	32(%rsp), %rdi
	movl	$21512, %esi
	call	cudaMalloc@PLT
	movl	$197, %esi
	movl	%eax, %edi
	call	_ZL11HandleError9cudaErrorPKci.constprop.0
	movq	16(%rsp), %rdi
	movq	%rbx, %rsi
	movl	$1, %ecx
	movabsq	$4294967303, %rbx
	movl	$21512, %edx
	call	cudaMemcpy@PLT
	movl	$200, %esi
	movl	%eax, %edi
	call	_ZL11HandleError9cudaErrorPKci.constprop.0
.L46:
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movq	%rbx, %rdx
	movl	$1, %ecx
	movq	%r13, %rdi
	movl	$1, %esi
	movq	%rbx, 52(%rsp)
	movl	$1, 60(%rsp)
	movq	%r13, 40(%rsp)
	movl	$1, 48(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L44
	movq	32(%rsp), %rsi
	movq	16(%rsp), %rdi
	call	_Z39__device_stub__Z17verschluessselungPlS_PlS_
.L44:
	call	cudaDeviceSynchronize@PLT
	movl	$211, %esi
	movl	%eax, %edi
	call	_ZL11HandleError9cudaErrorPKci.constprop.0
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movq	%rbx, %rdx
	movl	$1, %ecx
	movq	%r13, %rdi
	movl	$1, %esi
	movq	%rbx, 52(%rsp)
	movl	$1, 60(%rsp)
	movq	%r13, 40(%rsp)
	movl	$1, 48(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L45
	movq	24(%rsp), %rsi
	movq	32(%rsp), %rdi
	call	_Z39__device_stub__Z17entschluessselungPlS_PlS_
.L45:
	call	cudaDeviceSynchronize@PLT
	incl	%r15d
	movl	$217, %esi
	movl	%eax, %edi
	call	_ZL11HandleError9cudaErrorPKci.constprop.0
	cmpl	%r15d, %r12d
	jne	.L46
	movq	24(%rsp), %rsi
	movl	$2, %ecx
	movl	$21512, %edx
	leaq	21576(%rsp), %rbx
	movq	%rbx, %rdi
	call	cudaMemcpy@PLT
	movl	$221, %esi
	movl	%eax, %edi
	call	_ZL11HandleError9cudaErrorPKci.constprop.0
	movq	8(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movl	$224, %esi
	movl	%eax, %edi
	call	_ZL11HandleError9cudaErrorPKci.constprop.0
	movq	8(%rsp), %rdi
	call	cudaEventSynchronize@PLT
	movl	$225, %esi
	movl	%eax, %edi
	call	_ZL11HandleError9cudaErrorPKci.constprop.0
	movq	8(%rsp), %rdx
	movq	(%rsp), %rsi
	leaq	52(%rsp), %rdi
	call	cudaEventElapsedTime@PLT
	movl	$228, %esi
	movl	%eax, %edi
	call	_ZL11HandleError9cudaErrorPKci.constprop.0
	leaq	.LC11(%rip), %rsi
	movl	$2, %edi
	movb	$1, %al
	cvtss2sd	52(%rsp), %xmm0
	call	__printf_chk@PLT
	movq	%rbx, %rsi
	xorl	%ebx, %ebx
	movq	%rbp, %rdi
	call	_Z8unsplittPcPl
	leaq	.LC12(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
.L47:
	movsbl	0(%rbp,%rbx), %edi
	incq	%rbx
	call	putchar@PLT
	cmpq	$545, %rbx
	jne	.L47
	movq	%r14, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	(%rsp), %rdi
	call	cudaEventDestroy@PLT
	movl	$246, %esi
	movl	%eax, %edi
	call	_ZL11HandleError9cudaErrorPKci.constprop.0
	movq	8(%rsp), %rdi
	call	cudaEventDestroy@PLT
	movl	$247, %esi
	movl	%eax, %edi
	call	_ZL11HandleError9cudaErrorPKci.constprop.0
	movq	16(%rsp), %rdi
	call	cudaFree@PLT
	movl	$248, %esi
	movl	%eax, %edi
	call	_ZL11HandleError9cudaErrorPKci.constprop.0
	movq	24(%rsp), %rdi
	call	cudaFree@PLT
	movl	$249, %esi
	movl	%eax, %edi
	call	_ZL11HandleError9cudaErrorPKci.constprop.0
	movq	32(%rsp), %rdi
	call	cudaFree@PLT
	movl	$250, %esi
	movl	%eax, %edi
	call	_ZL11HandleError9cudaErrorPKci.constprop.0
	movq	48472(%rsp), %rax
	subq	%fs:40, %rax
	je	.L48
	call	__stack_chk_fail@PLT
.L48:
	addq	$48488, %rsp
	.cfi_def_cfa_offset 56
	xorl	%eax, %eax
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2030:
	.size	main, .-main
	.section	.rodata.str1.1
.LC13:
	.string	"_Z17entschluessselungPlS_"
.LC14:
	.string	"_Z17verschluessselungPlS_"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2060:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	pushq	$0
	.cfi_def_cfa_offset 24
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC13(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z17entschluessselungPlS_(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC14(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z17verschluessselungPlS_(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2060:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.rodata
	.type	CSWTCH.44, @object
	.size	CSWTCH.44, 4
CSWTCH.44:
	.byte	44
	.byte	46
	.byte	10
	.byte	32
	.section	.nvFatBinSegment,"aw"
	.align 8
	.type	_ZL15__fatDeviceText, @object
	.size	_ZL15__fatDeviceText, 24
_ZL15__fatDeviceText:
	.long	1180844977
	.long	1
	.quad	fatbinData
	.quad	0
	.local	_ZL20__cudaFatCubinHandle
	.comm	_ZL20__cudaFatCubinHandle,8,8
	.ident	"GCC: (Ubuntu 13.3.0-6ubuntu2~24.04) 13.3.0"
	.section	.note.GNU-stack,"",@progbits
	.section	.note.gnu.property,"a"
	.align 8
	.long	1f - 0f
	.long	4f - 1f
	.long	5
0:
	.string	"GNU"
1:
	.align 8
	.long	0xc0000002
	.long	3f - 2f
2:
	.long	0x3
3:
	.align 8
4:


#include <hip/hip_runtime.h>
#include <stdlib.h>
#include <stdio.h>

#define output 1

//Variablen
#define p 5
#define q 7
#define n 35
#define e 5
#define v 5
#define z 24
#define anzahl_Zeichen 2688
#define count_cores 384

/*
Klartext: K
Geheimtext: G
Verschluesselung: G = K^v mod n
Entschluesselung: K = G^e mod n
Index des CUDA Kerns: blockIdx.x blockIdx.y
Ein groesserer Text soll ver- und entschluesselt werden.
Dieser wird jedoch wie folgt veraendert: nur kleine Buchstaben, keine Sonderzeichen außer . und ,
Dafuer werden die einzellnen chars in Integer umgewandelt (<=30).
Somit ist eine Verarbeitung moeglich.
*/

static void HandleError( hipError_t err, const char *file, int line ) {
	if (err != hipSuccess) {
		printf( "%s in %s at line %d\n", hipGetErrorString( err ), file, line );
		exit( EXIT_FAILURE );
	}
}

#define HANDLE_ERROR( err ) (HandleError( err, __FILE__, __LINE__ ))

__global__ void verschluessselung(long int klartexte[], long int geheimtexte[])
{
	long int j, multi, x;
	
	//Fuer 384 Bloecke mit Threads
	long int threads = anzahl_Zeichen/count_cores;

	multi = x  = klartexte[threadIdx.x+blockIdx.x*threads];
	for (j = 1; j < v; j++)
		x *= multi;
	
	geheimtexte[threadIdx.x+blockIdx.x*threads] = x % n;
}

__global__ void entschluessselung(long int geheimtexte[], long int klartexte_pruefung[])
{
	long int j, multi, x;
	
	//Fuer 384 Bloecke mit Threads
	long int threads = anzahl_Zeichen/count_cores;
	
	multi = x  = geheimtexte[threadIdx.x+blockIdx.x*threads];
	for (j = 1; j < e; j++)
		x *= multi;
	
	klartexte_pruefung[threadIdx.x+blockIdx.x*threads] = x % n;
}

/*
 * Umwandlung von char in Zahl 0 bis 30 .
 */
void splitt(char text[], long int numbers[])
{
	int i;
	
	//Splitte Klartext
	for (i = 0; i < anzahl_Zeichen; i++)
	{
		long int number = (int)text[i];
		
		//char in int beginnend mit 0
		//Sonderzeichen
		if (number == 44)		//,
			number = 27;
		else if (number == 46)		//.
			number = 28;
		else if (number == 10)		//\n
			number = 29;
		else if (number == 32)		//' '
			number = 30;
		else				//a-z
			number -= 97;
		
		numbers[i] = number;
	}
}

/*
 * Umwandlung von Zahl 0 bis 30 zu char.
 */
void unsplitt(char text[], long int numbers[])
{
	int i;
	
	//Splitte Klartext
	for (i = 0; i < anzahl_Zeichen; i++)
	{
		long int number = numbers[i];
		char t;
		
		//int in char
		//Sonderzeichen
		if (number == 27)		//,
			t = ',';
		else if (number == 28)		//.
			number = t = '.';
		else if (number == 29)		//\n
			t = '\n';
		else if (number == 30)		//' '
			t = ' ';
		else				//a-z
			t = (char)(number+97);
			
		text[i] = t;
	}
}

int main(int argc, char *argv[]) {
	int i;
	hipEvent_t start, stop;
	float elapsedTime;
	int count_Threads;
	long int multi = 0;
	
	char klartext[anzahl_Zeichen+1];
	char klartext2[anzahl_Zeichen+1];
	long int kt_splitted[anzahl_Zeichen+1];
	long int kt_splitted2[anzahl_Zeichen+1];
	long int *dev_kt_splitted, *dev_kt_splitted2, *dev_gt_splitted;
	int size = sizeof(long int)*(anzahl_Zeichen+1);
	
	printf("\n-|| RSA mit CUDA ||-\n\n\n");
	
	//Klartetext erzeugen
	klartext[anzahl_Zeichen] = klartext2[anzahl_Zeichen] = '\0';
	strcpy(klartext, "hat der alte hexenmeister?sich doch einmal wegbegeben.?und nun sollen seine geister?auch nach meinem willen leben.?seine wort und werke?merkt ich und den brauch,?und mit geistesstaerke?tu ich wunder auch.?walle. walle?manche strecke,?dass, zum zwecke,?wasser fliesse?und mit reichem, vollem schwalle?zu dem bade sich ergiesse.?und nun komm, du alter besen.?nimm die schlechten lumpenhuellen ?bist schon lange knecht gewesen:?nun erfuelle meinen willen.?auf zwei beinen stehe,?oben sei ein kopf,?eile nun und gehe?mit dem wassertopf.?walle. walle?manche strecke,?dass, zum zwecke,?wasser fliesse?und mit reichem, vollem schwalle?zu dem bade sich ergiesse.?seht, er laeuft zum ufer nieder,?wahrlich. ist schon an dem flusse,?und mit blitzesschnelle wieder?ist er hier mit raschem gusse.?schon zum zweiten male.?wie das becken schwillt.?wie sich jede schale?voll mit wasser fuellt.?stehe. stehe.?denn wir haben?deiner gaben?vollgemessen. ?ach, ich merk es. wehe. wehe.?hab ich doch das wort vergessen.?ach, das wort, worauf am ende?er das wird, was er gewesen.?ach, er laeuft und bringt behende.?waerst du doch der alte besen.?immer neue guesse?bringt er schnell herein,?ach. und hundert fluesse?stuerzen auf mich ein.?nein, nicht laenger?kann ichs lassen ?will ihn fassen.?das ist tuecke.?ach. nun wird mir immer baenger.?welche mine. welche blicke.?o du ausgeburt der hoelle.?soll das ganze haus ersaufen??seh ich ueber jede schwelle?doch schon wasserstroeme laufen.?ein verruchter besen,?der nicht hoeren will.?stock, der du gewesen,?steh doch wieder still.?willst am ende?gar nicht lassen??will dich fassen,?will dich halten?und das alte holz behende?mit dem scharfen beile spalten.?seht da kommt er schleppend wieder.?wie ich mich nur auf dich werfe,?gleich, o kobold, liegst du nieder ?krachend trifft die glatte schaerfe.?wahrlich, brav getroffen.?seht, er ist entzwei.?und nun kann ich hoffen,?und ich atme frei.?wehe. wehe.?beide teile?stehn in eile?schon als knechte?voellig fertig in die hoehe.?helft mir, ach. ihr hohen maechte.?und sie laufen. nass und naesser?wirds im saal und auf den stufen.?welch entsetzliches gewaesser.?herr und meister. hoer mich rufen.  ?ach, da kommt der meister.?herr, die not ist gross.?die ich rief, die geister?werd ich nun nicht los.?in die ecke,?besen, besen.?seids gewesen.?denn als geister?ruft euch nur zu diesem zwecke,?erst hervor der alte meister.?                                                                                                                                                                                                                                                                                                    ");
	printf("\n\nDer Klartext ist %d Zeichen lang.\n", sizeof(klartext)/sizeof(char)-1);
	
	if (argc < 2)
	{
		printf("\nParameter fuer Groesse der Klartexted fehlt!\n");
		exit(0);
	}

	//lese Anzahl der Zeichen
	multi = atoi(argv[1]);

	if (multi < 1)
	{
		printf("\nAnzahl der Zeichen ist eine ganze positive Zahl.\n");
		exit(0);
	}
	
	printf("\nmulti: %d\n", multi);
	
	//Ausgabe
	if (output)
	{
		printf("\n\nAnfang des  Klartextes:\n\n");
		for (i = 0; i < anzahl_Zeichen; i++)
		{
			if (klartext[i] == '?')
				klartext[i] = '\n';
			if (i < 546)
				putchar(klartext[i]);
		}
		printf("\n\n");
	}
	
	//klartext2 mit a füllen
	for (i = 0; i < anzahl_Zeichen; i++)
	{
		klartext2[i] = 'a';
	}
	
	printf("Der Klartext wird nun verschluesselt und anschliessend entschluesselt.\n");
	
	//Chars in ints aufsplitten
	splitt(klartext, kt_splitted);
	
	//Variablen der Zeitmessung erstellen
	HANDLE_ERROR(hipEventCreate(&start));
	HANDLE_ERROR(hipEventCreate(&stop));
	
	//Start Zeitmessung
	HANDLE_ERROR(hipEventRecord(start, 0));
	
	//allokieren
	HANDLE_ERROR(hipMalloc((void **)&dev_kt_splitted, size));
	HANDLE_ERROR(hipMalloc((void **)&dev_kt_splitted2, size));
	HANDLE_ERROR(hipMalloc((void **)&dev_gt_splitted, size));
	
	//kopieren
	HANDLE_ERROR(hipMemcpy(dev_kt_splitted, kt_splitted, size, hipMemcpyHostToDevice));
	
	//Anzahl Threads pro Block
	count_Threads = anzahl_Zeichen/count_cores;
	
	for (i = 0; i < multi; i++)
	{
		//verschluesseln
		verschluessselung<<<count_cores, count_Threads>>>(dev_kt_splitted, dev_gt_splitted);
		
		//sync
		HANDLE_ERROR(hipDeviceSynchronize());
		
		//entschluesseln
		entschluessselung<<<count_cores, count_Threads>>>(dev_gt_splitted, dev_kt_splitted2);
		
		//sync
		HANDLE_ERROR(hipDeviceSynchronize());
	}
	
	//zurueckkopieren
	HANDLE_ERROR(hipMemcpy(kt_splitted2, dev_kt_splitted2, size, hipMemcpyDeviceToHost));
		
	//Ende der Zeitmessung
	HANDLE_ERROR(hipEventRecord(stop, 0));
	HANDLE_ERROR(hipEventSynchronize(stop));
	
	//Ausgabe der verstrichenen Zeit
	HANDLE_ERROR(hipEventElapsedTime(&elapsedTime, start, stop));
	printf("\nBeendet.\n\n\nverstrichene Zeit: %3.1f ms\n", elapsedTime);
	
	//ints wieder in char umwandeln
	unsplitt(klartext2, kt_splitted2);
	
	//Ausgabe
	if (output)
	{
		printf("\n\nDer Klartext lautet nun: (Anfang)\n\n");
		for (i = 0; i < 545; i++)
		{
			putchar(klartext2[i]);
		}
		printf("\n\n");
	}
	
	//freigeben
	HANDLE_ERROR(hipEventDestroy(start));
	HANDLE_ERROR(hipEventDestroy(stop));
	HANDLE_ERROR(hipFree(dev_kt_splitted));
	HANDLE_ERROR(hipFree(dev_kt_splitted2));
	HANDLE_ERROR(hipFree(dev_gt_splitted));
	return EXIT_SUCCESS;
}

	.text
	.file	"rsa.hip"
	.globl	_Z32__device_stub__verschluessselungPlS_ # -- Begin function _Z32__device_stub__verschluessselungPlS_
	.type	_Z32__device_stub__verschluessselungPlS_,@function
_Z32__device_stub__verschluessselungPlS_: # @_Z32__device_stub__verschluessselungPlS_
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$80, %rsp
	.cfi_def_cfa_offset 128
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	24(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	16(%rsp), %rcx
	movq	%rsi, (%rcx)
	leaq	64(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%rbx)
	leaq	48(%rsp), %r14
	leaq	32(%rsp), %r15
	leaq	8(%rsp), %r12
	movq	%rsp, %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z17verschluessselungPlS_, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$96, %rsp
	.cfi_adjust_cfa_offset -96
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end0:
	.size	_Z32__device_stub__verschluessselungPlS_, .Lfunc_end0-_Z32__device_stub__verschluessselungPlS_
	.cfi_endproc
                                        # -- End function
	.globl	_Z32__device_stub__entschluessselungPlS_ # -- Begin function _Z32__device_stub__entschluessselungPlS_
	.type	_Z32__device_stub__entschluessselungPlS_,@function
_Z32__device_stub__entschluessselungPlS_: # @_Z32__device_stub__entschluessselungPlS_
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$80, %rsp
	.cfi_def_cfa_offset 128
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	24(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	16(%rsp), %rcx
	movq	%rsi, (%rcx)
	leaq	64(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rcx, 8(%rbx)
	leaq	48(%rsp), %r14
	leaq	32(%rsp), %r15
	leaq	8(%rsp), %r12
	movq	%rsp, %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z17entschluessselungPlS_, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$96, %rsp
	.cfi_adjust_cfa_offset -96
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end1:
	.size	_Z32__device_stub__entschluessselungPlS_, .Lfunc_end1-_Z32__device_stub__entschluessselungPlS_
	.cfi_endproc
                                        # -- End function
	.globl	_Z6splittPcPl                   # -- Begin function _Z6splittPcPl
	.type	_Z6splittPcPl,@function
_Z6splittPcPl:                          # @_Z6splittPcPl
	.cfi_startproc
# %bb.0:
	xorl	%eax, %eax
.LBB2_1:                                # =>This Inner Loop Header: Depth=1
	movzbl	(%rdi,%rax), %ecx
	cmpl	$43, %ecx
	jg	.LBB2_5
# %bb.2:                                #   in Loop: Header=BB2_1 Depth=1
	cmpl	$10, %ecx
	je	.LBB2_9
# %bb.3:                                #   in Loop: Header=BB2_1 Depth=1
	cmpl	$32, %ecx
	jne	.LBB2_10
# %bb.4:                                #   in Loop: Header=BB2_1 Depth=1
	movl	$30, %ecx
	jmp	.LBB2_11
.LBB2_5:                                #   in Loop: Header=BB2_1 Depth=1
	cmpl	$44, %ecx
	je	.LBB2_6
# %bb.7:                                #   in Loop: Header=BB2_1 Depth=1
	cmpl	$46, %ecx
	jne	.LBB2_10
# %bb.8:                                #   in Loop: Header=BB2_1 Depth=1
	movl	$28, %ecx
	jmp	.LBB2_11
.LBB2_10:                               #   in Loop: Header=BB2_1 Depth=1
	movsbq	%cl, %rcx
	addq	$-97, %rcx
	jmp	.LBB2_11
.LBB2_9:                                #   in Loop: Header=BB2_1 Depth=1
	movl	$29, %ecx
	jmp	.LBB2_11
.LBB2_6:                                #   in Loop: Header=BB2_1 Depth=1
	movl	$27, %ecx
.LBB2_11:                               #   in Loop: Header=BB2_1 Depth=1
	movq	%rcx, (%rsi,%rax,8)
	incq	%rax
	cmpq	$2688, %rax                     # imm = 0xA80
	jne	.LBB2_1
# %bb.12:
	retq
.Lfunc_end2:
	.size	_Z6splittPcPl, .Lfunc_end2-_Z6splittPcPl
	.cfi_endproc
                                        # -- End function
	.globl	_Z8unsplittPcPl                 # -- Begin function _Z8unsplittPcPl
	.type	_Z8unsplittPcPl,@function
_Z8unsplittPcPl:                        # @_Z8unsplittPcPl
	.cfi_startproc
# %bb.0:
	xorl	%eax, %eax
.LBB3_1:                                # =>This Inner Loop Header: Depth=1
	movq	(%rsi,%rax,8), %rdx
	leaq	-27(%rdx), %r8
	addb	$97, %dl
	leal	(,%r8,8), %ecx
	movl	$537538092, %r9d                # imm = 0x200A2E2C
                                        # kill: def $cl killed $cl killed $ecx
	shrl	%cl, %r9d
	cmpq	$4, %r8
	movzbl	%dl, %ecx
	cmovbl	%r9d, %ecx
	movb	%cl, (%rdi,%rax)
	incq	%rax
	cmpq	$2688, %rax                     # imm = 0xA80
	jne	.LBB3_1
# %bb.2:
	retq
.Lfunc_end3:
	.size	_Z8unsplittPcPl, .Lfunc_end3-_Z8unsplittPcPl
	.cfi_endproc
                                        # -- End function
	.globl	main                            # -- Begin function main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	subq	$48488, %rsp                    # imm = 0xBD68
	.cfi_def_cfa_offset 48528
	.cfi_offset %rbx, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movq	%rsi, %rbx
	movl	%edi, %ebp
	movl	$.Lstr, %edi
	callq	puts@PLT
	movb	$0, 2736(%rsp)
	leaq	2752(%rsp), %rdi
	movl	$.L.str.1, %esi
	movl	$2689, %edx                     # imm = 0xA81
	callq	memcpy@PLT
	movl	$.L.str.2, %edi
	movl	$2688, %esi                     # imm = 0xA80
	xorl	%eax, %eax
	callq	printf
	cmpl	$1, %ebp
	jle	.LBB4_1
# %bb.3:
	movq	8(%rbx), %rdi
	callq	atoi
	testl	%eax, %eax
	jle	.LBB4_4
# %bb.5:
	movl	%eax, %ebx
	movl	%eax, %esi
	xorl	%r14d, %r14d
	movl	$.L.str.5, %edi
	xorl	%eax, %eax
	callq	printf
	movl	$.Lstr.1, %edi
	callq	puts@PLT
	movl	$10, %ebp
.LBB4_6:                                # =>This Inner Loop Header: Depth=1
	movzbl	2752(%rsp,%r14), %eax
	cmpb	$63, %al
	cmovel	%ebp, %eax
	movb	%al, 2752(%rsp,%r14)
	cmpq	$545, %r14                      # imm = 0x221
	ja	.LBB4_8
# %bb.7:                                #   in Loop: Header=BB4_6 Depth=1
	movsbl	%al, %edi
	callq	putchar
.LBB4_8:                                #   in Loop: Header=BB4_6 Depth=1
	incq	%r14
	cmpq	$2688, %r14                     # imm = 0xA80
	jne	.LBB4_6
# %bb.9:
	movl	$.Lstr.5, %edi
	callq	puts@PLT
	leaq	48(%rsp), %rdi
	movl	$2688, %edx                     # imm = 0xA80
	movl	$97, %esi
	callq	memset@PLT
	movl	$.Lstr.3, %edi
	callq	puts@PLT
	xorl	%eax, %eax
.LBB4_10:                               # =>This Inner Loop Header: Depth=1
	movzbl	2752(%rsp,%rax), %ecx
	cmpl	$43, %ecx
	jg	.LBB4_14
# %bb.11:                               #   in Loop: Header=BB4_10 Depth=1
	cmpl	$10, %ecx
	je	.LBB4_18
# %bb.12:                               #   in Loop: Header=BB4_10 Depth=1
	cmpl	$32, %ecx
	jne	.LBB4_19
# %bb.13:                               #   in Loop: Header=BB4_10 Depth=1
	movl	$30, %ecx
	jmp	.LBB4_20
.LBB4_14:                               #   in Loop: Header=BB4_10 Depth=1
	cmpl	$44, %ecx
	je	.LBB4_15
# %bb.16:                               #   in Loop: Header=BB4_10 Depth=1
	cmpl	$46, %ecx
	jne	.LBB4_19
# %bb.17:                               #   in Loop: Header=BB4_10 Depth=1
	movl	$28, %ecx
	jmp	.LBB4_20
.LBB4_19:                               #   in Loop: Header=BB4_10 Depth=1
	movsbq	%cl, %rcx
	addq	$-97, %rcx
	jmp	.LBB4_20
.LBB4_18:                               #   in Loop: Header=BB4_10 Depth=1
	movl	$29, %ecx
	jmp	.LBB4_20
.LBB4_15:                               #   in Loop: Header=BB4_10 Depth=1
	movl	$27, %ecx
.LBB4_20:                               #   in Loop: Header=BB4_10 Depth=1
	movq	%rcx, 26976(%rsp,%rax,8)
	incq	%rax
	cmpq	$2688, %rax                     # imm = 0xA80
	jne	.LBB4_10
# %bb.21:                               # %.lr.ph.preheader
	movabsq	$4294967303, %r14               # imm = 0x100000007
	leaq	32(%rsp), %r15
	movq	%r15, %rdi
	callq	hipEventCreate
	movl	%eax, %edi
	movl	$190, %esi
	callq	_ZL11HandleError10hipError_tPKci
	movq	%rsp, %rdi
	callq	hipEventCreate
	movl	%eax, %edi
	movl	$191, %esi
	callq	_ZL11HandleError10hipError_tPKci
	movq	(%r15), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movl	%eax, %edi
	movl	$194, %esi
	callq	_ZL11HandleError10hipError_tPKci
	leaq	24(%rsp), %r15
	movl	$21512, %esi                    # imm = 0x5408
	movq	%r15, %rdi
	callq	hipMalloc
	movl	%eax, %edi
	movl	$197, %esi
	callq	_ZL11HandleError10hipError_tPKci
	leaq	16(%rsp), %rdi
	movl	$21512, %esi                    # imm = 0x5408
	callq	hipMalloc
	movl	%eax, %edi
	movl	$198, %esi
	callq	_ZL11HandleError10hipError_tPKci
	leaq	8(%rsp), %rdi
	movl	$21512, %esi                    # imm = 0x5408
	callq	hipMalloc
	movl	%eax, %edi
	movl	$199, %esi
	callq	_ZL11HandleError10hipError_tPKci
	movq	(%r15), %rdi
	leaq	26976(%rsp), %rsi
	movl	$21512, %edx                    # imm = 0x5408
	movl	$1, %ecx
	callq	hipMemcpy
	movl	%eax, %edi
	movl	$202, %esi
	callq	_ZL11HandleError10hipError_tPKci
	leaq	377(%r14), %r15
.LBB4_22:                               # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movq	%r15, %rdi
	movl	$1, %esi
	movq	%r14, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB4_24
# %bb.23:                               #   in Loop: Header=BB4_22 Depth=1
	movq	24(%rsp), %rdi
	movq	8(%rsp), %rsi
	callq	_Z32__device_stub__verschluessselungPlS_
.LBB4_24:                               #   in Loop: Header=BB4_22 Depth=1
	callq	hipDeviceSynchronize
	movl	%eax, %edi
	movl	$213, %esi
	callq	_ZL11HandleError10hipError_tPKci
	movq	%r15, %rdi
	movl	$1, %esi
	movq	%r14, %rdx
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB4_26
# %bb.25:                               #   in Loop: Header=BB4_22 Depth=1
	movq	8(%rsp), %rdi
	movq	16(%rsp), %rsi
	callq	_Z32__device_stub__entschluessselungPlS_
.LBB4_26:                               #   in Loop: Header=BB4_22 Depth=1
	callq	hipDeviceSynchronize
	movl	%eax, %edi
	movl	$219, %esi
	callq	_ZL11HandleError10hipError_tPKci
	decl	%ebx
	jne	.LBB4_22
# %bb.27:                               # %._crit_edge
	movq	16(%rsp), %rsi
	leaq	5456(%rsp), %rdi
	movl	$21512, %edx                    # imm = 0x5408
	movl	$2, %ecx
	callq	hipMemcpy
	movl	%eax, %edi
	movl	$223, %esi
	callq	_ZL11HandleError10hipError_tPKci
	movq	(%rsp), %rdi
	xorl	%r14d, %r14d
	xorl	%esi, %esi
	callq	hipEventRecord
	movl	%eax, %edi
	movl	$226, %esi
	callq	_ZL11HandleError10hipError_tPKci
	movq	(%rsp), %rdi
	callq	hipEventSynchronize
	movl	%eax, %edi
	movl	$227, %esi
	callq	_ZL11HandleError10hipError_tPKci
	movq	32(%rsp), %rsi
	movq	(%rsp), %rdx
	leaq	44(%rsp), %rbx
	movq	%rbx, %rdi
	callq	hipEventElapsedTime
	movl	%eax, %edi
	movl	$230, %esi
	callq	_ZL11HandleError10hipError_tPKci
	cvtss2sd	(%rbx), %xmm0
	movl	$.L.str.10, %edi
	movb	$1, %al
	callq	printf
.LBB4_28:                               # =>This Inner Loop Header: Depth=1
	movq	5456(%rsp,%r14,8), %rax
	leaq	-27(%rax), %rdx
	addb	$97, %al
	leal	(,%rdx,8), %ecx
	movl	$537538092, %esi                # imm = 0x200A2E2C
                                        # kill: def $cl killed $cl killed $ecx
	shrl	%cl, %esi
	cmpq	$4, %rdx
	movzbl	%al, %eax
	cmovbl	%esi, %eax
	movb	%al, 48(%rsp,%r14)
	incq	%r14
	cmpq	$2688, %r14                     # imm = 0xA80
	jne	.LBB4_28
# %bb.29:                               # %_Z8unsplittPcPl.exit
	movl	$.Lstr.4, %edi
	callq	puts@PLT
	xorl	%ebx, %ebx
.LBB4_30:                               # =>This Inner Loop Header: Depth=1
	movsbl	48(%rsp,%rbx), %edi
	callq	putchar
	incq	%rbx
	cmpq	$545, %rbx                      # imm = 0x221
	jne	.LBB4_30
# %bb.31:
	movl	$.Lstr.5, %edi
	callq	puts@PLT
	movq	32(%rsp), %rdi
	callq	hipEventDestroy
	movl	%eax, %edi
	movl	$248, %esi
	callq	_ZL11HandleError10hipError_tPKci
	movq	(%rsp), %rdi
	callq	hipEventDestroy
	movl	%eax, %edi
	movl	$249, %esi
	callq	_ZL11HandleError10hipError_tPKci
	movq	24(%rsp), %rdi
	callq	hipFree
	movl	%eax, %edi
	movl	$250, %esi
	callq	_ZL11HandleError10hipError_tPKci
	movq	16(%rsp), %rdi
	callq	hipFree
	movl	%eax, %edi
	movl	$251, %esi
	callq	_ZL11HandleError10hipError_tPKci
	movq	8(%rsp), %rdi
	callq	hipFree
	movl	%eax, %edi
	movl	$252, %esi
	callq	_ZL11HandleError10hipError_tPKci
	xorl	%eax, %eax
	addq	$48488, %rsp                    # imm = 0xBD68
	.cfi_def_cfa_offset 40
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.LBB4_1:
	.cfi_def_cfa_offset 48528
	movl	$.Lstr.7, %edi
	jmp	.LBB4_2
.LBB4_4:
	movl	$.Lstr.6, %edi
.LBB4_2:
	callq	puts@PLT
	xorl	%edi, %edi
	callq	exit
.Lfunc_end4:
	.size	main, .Lfunc_end4-main
	.cfi_endproc
                                        # -- End function
	.type	_ZL11HandleError10hipError_tPKci,@function # -- Begin function _ZL11HandleError10hipError_tPKci
_ZL11HandleError10hipError_tPKci:       # @_ZL11HandleError10hipError_tPKci
	.cfi_startproc
# %bb.0:
	testl	%edi, %edi
	jne	.LBB5_2
# %bb.1:
	retq
.LBB5_2:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset %rbx, -16
	movl	%esi, %ebx
	callq	hipGetErrorString
	movl	$.L.str.12, %edi
	movl	$.L.str.9, %edx
	movq	%rax, %rsi
	movl	%ebx, %ecx
	xorl	%eax, %eax
	callq	printf
	movl	$1, %edi
	callq	exit
.Lfunc_end5:
	.size	_ZL11HandleError10hipError_tPKci, .Lfunc_end5-_ZL11HandleError10hipError_tPKci
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_ctor
	.type	__hip_module_ctor,@function
__hip_module_ctor:                      # @__hip_module_ctor
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	subq	$32, %rsp
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -16
	movq	__hip_gpubin_handle(%rip), %rbx
	testq	%rbx, %rbx
	jne	.LBB6_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rbx
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB6_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z17verschluessselungPlS_, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z17entschluessselungPlS_, %esi
	movl	$.L__unnamed_2, %edx
	movl	$.L__unnamed_2, %ecx
	movq	%rbx, %rdi
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end6:
	.size	__hip_module_ctor, .Lfunc_end6-__hip_module_ctor
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_dtor
	.type	__hip_module_dtor,@function
__hip_module_dtor:                      # @__hip_module_dtor
	.cfi_startproc
# %bb.0:
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	je	.LBB7_2
# %bb.1:
	pushq	%rax
	.cfi_def_cfa_offset 16
	callq	__hipUnregisterFatBinary
	movq	$0, __hip_gpubin_handle(%rip)
	addq	$8, %rsp
	.cfi_def_cfa_offset 8
.LBB7_2:
	retq
.Lfunc_end7:
	.size	__hip_module_dtor, .Lfunc_end7-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z17verschluessselungPlS_,@object # @_Z17verschluessselungPlS_
	.section	.rodata,"a",@progbits
	.globl	_Z17verschluessselungPlS_
	.p2align	3, 0x0
_Z17verschluessselungPlS_:
	.quad	_Z32__device_stub__verschluessselungPlS_
	.size	_Z17verschluessselungPlS_, 8

	.type	_Z17entschluessselungPlS_,@object # @_Z17entschluessselungPlS_
	.globl	_Z17entschluessselungPlS_
	.p2align	3, 0x0
_Z17entschluessselungPlS_:
	.quad	_Z32__device_stub__entschluessselungPlS_
	.size	_Z17entschluessselungPlS_, 8

	.type	.L.str.1,@object                # @.str.1
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str.1:
	.asciz	"hat der alte hexenmeister?sich doch einmal wegbegeben.?und nun sollen seine geister?auch nach meinem willen leben.?seine wort und werke?merkt ich und den brauch,?und mit geistesstaerke?tu ich wunder auch.?walle. walle?manche strecke,?dass, zum zwecke,?wasser fliesse?und mit reichem, vollem schwalle?zu dem bade sich ergiesse.?und nun komm, du alter besen.?nimm die schlechten lumpenhuellen ?bist schon lange knecht gewesen:?nun erfuelle meinen willen.?auf zwei beinen stehe,?oben sei ein kopf,?eile nun und gehe?mit dem wassertopf.?walle. walle?manche strecke,?dass, zum zwecke,?wasser fliesse?und mit reichem, vollem schwalle?zu dem bade sich ergiesse.?seht, er laeuft zum ufer nieder,?wahrlich. ist schon an dem flusse,?und mit blitzesschnelle wieder?ist er hier mit raschem gusse.?schon zum zweiten male.?wie das becken schwillt.?wie sich jede schale?voll mit wasser fuellt.?stehe. stehe.?denn wir haben?deiner gaben?vollgemessen. ?ach, ich merk es. wehe. wehe.?hab ich doch das wort vergessen.?ach, das wort, worauf am ende?er das wird, was er gewesen.?ach, er laeuft und bringt behende.?waerst du doch der alte besen.?immer neue guesse?bringt er schnell herein,?ach. und hundert fluesse?stuerzen auf mich ein.?nein, nicht laenger?kann ichs lassen ?will ihn fassen.?das ist tuecke.?ach. nun wird mir immer baenger.?welche mine. welche blicke.?o du ausgeburt der hoelle.?soll das ganze haus ersaufen??seh ich ueber jede schwelle?doch schon wasserstroeme laufen.?ein verruchter besen,?der nicht hoeren will.?stock, der du gewesen,?steh doch wieder still.?willst am ende?gar nicht lassen??will dich fassen,?will dich halten?und das alte holz behende?mit dem scharfen beile spalten.?seht da kommt er schleppend wieder.?wie ich mich nur auf dich werfe,?gleich, o kobold, liegst du nieder ?krachend trifft die glatte schaerfe.?wahrlich, brav getroffen.?seht, er ist entzwei.?und nun kann ich hoffen,?und ich atme frei.?wehe. wehe.?beide teile?stehn in eile?schon als knechte?voellig fertig in die hoehe.?helft mir, ach. ihr hohen maechte.?und sie laufen. nass und naesser?wirds im saal und auf den stufen.?welch entsetzliches gewaesser.?herr und meister. hoer mich rufen.  ?ach, da kommt der meister.?herr, die not ist gross.?die ich rief, die geister?werd ich nun nicht los.?in die ecke,?besen, besen.?seids gewesen.?denn als geister?ruft euch nur zu diesem zwecke,?erst hervor der alte meister.?                                                                                                                                                                                                                                                                                                    "
	.size	.L.str.1, 2689

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	"\n\nDer Klartext ist %d Zeichen lang.\n"
	.size	.L.str.2, 37

	.type	.L.str.5,@object                # @.str.5
.L.str.5:
	.asciz	"\nmulti: %d\n"
	.size	.L.str.5, 12

	.type	.L.str.9,@object                # @.str.9
.L.str.9:
	.asciz	"/home/ubuntu/Datasets/stackv2/train-structured-repos-hip-sm89/TBoonX/krypto_cuda/master/rsa.hip"
	.size	.L.str.9, 96

	.type	.L.str.10,@object               # @.str.10
.L.str.10:
	.asciz	"\nBeendet.\n\n\nverstrichene Zeit: %3.1f ms\n"
	.size	.L.str.10, 41

	.type	.L.str.12,@object               # @.str.12
.L.str.12:
	.asciz	"%s in %s at line %d\n"
	.size	.L.str.12, 21

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z17verschluessselungPlS_"
	.size	.L__unnamed_1, 26

	.type	.L__unnamed_2,@object           # @1
.L__unnamed_2:
	.asciz	"_Z17entschluessselungPlS_"
	.size	.L__unnamed_2, 26

	.type	__hip_fatbin_wrapper,@object    # @__hip_fatbin_wrapper
	.section	.hipFatBinSegment,"a",@progbits
	.p2align	3, 0x0
__hip_fatbin_wrapper:
	.long	1212764230                      # 0x48495046
	.long	1                               # 0x1
	.quad	__hip_fatbin
	.quad	0
	.size	__hip_fatbin_wrapper, 24

	.type	__hip_gpubin_handle,@object     # @__hip_gpubin_handle
	.local	__hip_gpubin_handle
	.comm	__hip_gpubin_handle,8,8
	.section	.init_array,"aw",@init_array
	.p2align	3, 0x0
	.quad	__hip_module_ctor
	.type	__hip_cuid_,@object             # @__hip_cuid_
	.bss
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               # 0x0
	.size	__hip_cuid_, 1

	.type	.Lstr,@object                   # @str
	.section	.rodata.str1.1,"aMS",@progbits,1
.Lstr:
	.asciz	"\n-|| RSA mit CUDA ||-\n\n"
	.size	.Lstr, 24

	.type	.Lstr.1,@object                 # @str.1
.Lstr.1:
	.asciz	"\n\nAnfang des  Klartextes:\n"
	.size	.Lstr.1, 27

	.type	.Lstr.3,@object                 # @str.3
.Lstr.3:
	.asciz	"Der Klartext wird nun verschluesselt und anschliessend entschluesselt."
	.size	.Lstr.3, 71

	.type	.Lstr.4,@object                 # @str.4
.Lstr.4:
	.asciz	"\n\nDer Klartext lautet nun: (Anfang)\n"
	.size	.Lstr.4, 37

	.type	.Lstr.5,@object                 # @str.5
.Lstr.5:
	.asciz	"\n"
	.size	.Lstr.5, 2

	.type	.Lstr.6,@object                 # @str.6
.Lstr.6:
	.asciz	"\nAnzahl der Zeichen ist eine ganze positive Zahl."
	.size	.Lstr.6, 50

	.type	.Lstr.7,@object                 # @str.7
.Lstr.7:
	.asciz	"\nParameter fuer Groesse der Klartexted fehlt!"
	.size	.Lstr.7, 46

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym _Z32__device_stub__verschluessselungPlS_
	.addrsig_sym _Z32__device_stub__entschluessselungPlS_
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z17verschluessselungPlS_
	.addrsig_sym _Z17entschluessselungPlS_
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z17verschluessselungPlS_ ; -- Begin function _Z17verschluessselungPlS_
	.globl	_Z17verschluessselungPlS_
	.p2align	8
	.type	_Z17verschluessselungPlS_,@function
_Z17verschluessselungPlS_:              ; @_Z17verschluessselungPlS_
; %bb.0:
	s_load_b128 s[0:3], s[0:1], 0x0
	s_mul_hi_u32 s5, s15, 7
	s_mul_i32 s4, s15, 7
	v_lshlrev_b32_e32 v1, 3, v0
	s_lshl_b64 s[4:5], s[4:5], 3
	s_waitcnt lgkmcnt(0)
	s_add_u32 s0, s0, s4
	s_addc_u32 s1, s1, s5
	global_load_b64 v[3:4], v1, s[0:1]
	s_mov_b64 s[0:1], 4
	s_waitcnt vmcnt(0)
	v_dual_mov_b32 v6, v4 :: v_dual_mov_b32 v5, v3
.LBB0_1:                                ; =>This Inner Loop Header: Depth=1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_lo_u32 v6, v6, v3
	v_mul_lo_u32 v7, v5, v4
	v_mad_u64_u32 v[1:2], null, v5, v3, 0
	s_add_u32 s0, s0, -1
	s_addc_u32 s1, s1, -1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	s_cmp_lg_u64 s[0:1], 0
	v_add3_u32 v2, v2, v7, v6
	s_delay_alu instid0(VALU_DEP_1)
	v_dual_mov_b32 v6, v2 :: v_dual_mov_b32 v5, v1
	s_cbranch_scc1 .LBB0_1
; %bb.2:
	s_add_u32 s0, 0, 0x50748140
	s_addc_u32 s1, 0, 31
	s_mul_hi_u32 s7, s0, 0xffffffdd
	s_add_i32 s1, s1, 0x75074e8
	s_mul_i32 s6, s0, 0xffffffdd
	s_mul_i32 s8, s1, 0xffffffdd
	s_mul_hi_u32 s9, s0, s6
	s_add_i32 s7, s7, s8
	v_ashrrev_i32_e32 v5, 31, v2
	s_sub_i32 s7, s7, s0
	s_mul_hi_u32 s8, s1, s6
	s_mul_i32 s10, s0, s7
	s_mul_hi_u32 s11, s0, s7
	s_mul_i32 s6, s1, s6
	s_add_u32 s9, s9, s10
	s_addc_u32 s10, 0, s11
	s_mul_hi_u32 s11, s1, s7
	s_add_u32 s6, s9, s6
	s_addc_u32 s6, s10, s8
	s_mul_i32 s7, s1, s7
	v_add_co_u32 v1, vcc_lo, v1, v5
	s_addc_u32 s8, s11, 0
	s_add_u32 s6, s6, s7
	s_addc_u32 s7, 0, s8
	v_add_co_u32 v6, s0, s0, s6
	v_add_co_ci_u32_e32 v3, vcc_lo, v2, v5, vcc_lo
	v_xor_b32_e32 v7, v1, v5
	s_cmp_lg_u32 s0, 0
	v_lshlrev_b32_e32 v0, 3, v0
	s_addc_u32 s0, s1, s7
	v_xor_b32_e32 v9, v3, v5
	v_mad_u64_u32 v[1:2], null, v7, s0, 0
	v_mul_hi_u32 v8, v7, v6
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mad_u64_u32 v[3:4], null, v9, v6, 0
	v_add_co_u32 v6, vcc_lo, v8, v1
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_add_co_ci_u32_e32 v8, vcc_lo, 0, v2, vcc_lo
	v_mad_u64_u32 v[1:2], null, v9, s0, 0
	v_add_co_u32 v3, vcc_lo, v6, v3
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_ci_u32_e32 v3, vcc_lo, v8, v4, vcc_lo
	v_add_co_ci_u32_e32 v2, vcc_lo, 0, v2, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v3, vcc_lo, v3, v1
	v_add_co_ci_u32_e32 v6, vcc_lo, 0, v2, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[1:2], null, v3, 35, 0
	v_mad_u64_u32 v[3:4], null, v6, 35, v[2:3]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_sub_co_u32 v1, vcc_lo, v7, v1
	v_sub_co_ci_u32_e32 v2, vcc_lo, v9, v3, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_sub_co_u32 v3, vcc_lo, v1, 35
	v_subrev_co_ci_u32_e32 v4, vcc_lo, 0, v2, vcc_lo
	s_delay_alu instid0(VALU_DEP_2)
	v_cmp_lt_u32_e32 vcc_lo, 34, v3
	v_cmp_eq_u32_e64 s0, 0, v2
	v_cndmask_b32_e64 v6, 0, -1, vcc_lo
	v_cmp_lt_u32_e32 vcc_lo, 34, v1
	v_cndmask_b32_e64 v7, 0, -1, vcc_lo
	v_cmp_eq_u32_e32 vcc_lo, 0, v4
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_cndmask_b32_e32 v6, -1, v6, vcc_lo
	v_sub_co_u32 v8, vcc_lo, v3, 35
	v_subrev_co_ci_u32_e32 v9, vcc_lo, 0, v4, vcc_lo
	v_cmp_ne_u32_e32 vcc_lo, 0, v6
	v_cndmask_b32_e64 v6, -1, v7, s0
	s_add_u32 s0, s2, s4
	s_addc_u32 s1, s3, s5
	v_dual_cndmask_b32 v3, v3, v8 :: v_dual_cndmask_b32 v4, v4, v9
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cmp_ne_u32_e32 vcc_lo, 0, v6
	v_dual_cndmask_b32 v1, v1, v3 :: v_dual_cndmask_b32 v2, v2, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_xor_b32_e32 v1, v1, v5
	v_xor_b32_e32 v2, v2, v5
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_sub_co_u32 v1, vcc_lo, v1, v5
	v_sub_co_ci_u32_e32 v2, vcc_lo, v2, v5, vcc_lo
	global_store_b64 v0, v[1:2], s[0:1]
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z17verschluessselungPlS_
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 16
		.amdhsa_user_sgpr_count 15
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 10
		.amdhsa_next_free_sgpr 16
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end0:
	.size	_Z17verschluessselungPlS_, .Lfunc_end0-_Z17verschluessselungPlS_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 572
; NumSgprs: 18
; NumVgprs: 10
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 10
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.protected	_Z17entschluessselungPlS_ ; -- Begin function _Z17entschluessselungPlS_
	.globl	_Z17entschluessselungPlS_
	.p2align	8
	.type	_Z17entschluessselungPlS_,@function
_Z17entschluessselungPlS_:              ; @_Z17entschluessselungPlS_
; %bb.0:
	s_load_b128 s[0:3], s[0:1], 0x0
	s_mul_hi_u32 s5, s15, 7
	s_mul_i32 s4, s15, 7
	v_lshlrev_b32_e32 v1, 3, v0
	s_lshl_b64 s[4:5], s[4:5], 3
	s_waitcnt lgkmcnt(0)
	s_add_u32 s0, s0, s4
	s_addc_u32 s1, s1, s5
	global_load_b64 v[3:4], v1, s[0:1]
	s_mov_b64 s[0:1], 4
	s_waitcnt vmcnt(0)
	v_dual_mov_b32 v6, v4 :: v_dual_mov_b32 v5, v3
.LBB1_1:                                ; =>This Inner Loop Header: Depth=1
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mul_lo_u32 v6, v6, v3
	v_mul_lo_u32 v7, v5, v4
	v_mad_u64_u32 v[1:2], null, v5, v3, 0
	s_add_u32 s0, s0, -1
	s_addc_u32 s1, s1, -1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	s_cmp_lg_u64 s[0:1], 0
	v_add3_u32 v2, v2, v7, v6
	s_delay_alu instid0(VALU_DEP_1)
	v_dual_mov_b32 v6, v2 :: v_dual_mov_b32 v5, v1
	s_cbranch_scc1 .LBB1_1
; %bb.2:
	s_add_u32 s0, 0, 0x50748140
	s_addc_u32 s1, 0, 31
	s_mul_hi_u32 s7, s0, 0xffffffdd
	s_add_i32 s1, s1, 0x75074e8
	s_mul_i32 s6, s0, 0xffffffdd
	s_mul_i32 s8, s1, 0xffffffdd
	s_mul_hi_u32 s9, s0, s6
	s_add_i32 s7, s7, s8
	v_ashrrev_i32_e32 v5, 31, v2
	s_sub_i32 s7, s7, s0
	s_mul_hi_u32 s8, s1, s6
	s_mul_i32 s10, s0, s7
	s_mul_hi_u32 s11, s0, s7
	s_mul_i32 s6, s1, s6
	s_add_u32 s9, s9, s10
	s_addc_u32 s10, 0, s11
	s_mul_hi_u32 s11, s1, s7
	s_add_u32 s6, s9, s6
	s_addc_u32 s6, s10, s8
	s_mul_i32 s7, s1, s7
	v_add_co_u32 v1, vcc_lo, v1, v5
	s_addc_u32 s8, s11, 0
	s_add_u32 s6, s6, s7
	s_addc_u32 s7, 0, s8
	v_add_co_u32 v6, s0, s0, s6
	v_add_co_ci_u32_e32 v3, vcc_lo, v2, v5, vcc_lo
	v_xor_b32_e32 v7, v1, v5
	s_cmp_lg_u32 s0, 0
	v_lshlrev_b32_e32 v0, 3, v0
	s_addc_u32 s0, s1, s7
	v_xor_b32_e32 v9, v3, v5
	v_mad_u64_u32 v[1:2], null, v7, s0, 0
	v_mul_hi_u32 v8, v7, v6
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_mad_u64_u32 v[3:4], null, v9, v6, 0
	v_add_co_u32 v6, vcc_lo, v8, v1
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_1) | instid1(VALU_DEP_3)
	v_add_co_ci_u32_e32 v8, vcc_lo, 0, v2, vcc_lo
	v_mad_u64_u32 v[1:2], null, v9, s0, 0
	v_add_co_u32 v3, vcc_lo, v6, v3
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_add_co_ci_u32_e32 v3, vcc_lo, v8, v4, vcc_lo
	v_add_co_ci_u32_e32 v2, vcc_lo, 0, v2, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v3, vcc_lo, v3, v1
	v_add_co_ci_u32_e32 v6, vcc_lo, 0, v2, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[1:2], null, v3, 35, 0
	v_mad_u64_u32 v[3:4], null, v6, 35, v[2:3]
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_sub_co_u32 v1, vcc_lo, v7, v1
	v_sub_co_ci_u32_e32 v2, vcc_lo, v9, v3, vcc_lo
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_sub_co_u32 v3, vcc_lo, v1, 35
	v_subrev_co_ci_u32_e32 v4, vcc_lo, 0, v2, vcc_lo
	s_delay_alu instid0(VALU_DEP_2)
	v_cmp_lt_u32_e32 vcc_lo, 34, v3
	v_cmp_eq_u32_e64 s0, 0, v2
	v_cndmask_b32_e64 v6, 0, -1, vcc_lo
	v_cmp_lt_u32_e32 vcc_lo, 34, v1
	v_cndmask_b32_e64 v7, 0, -1, vcc_lo
	v_cmp_eq_u32_e32 vcc_lo, 0, v4
	s_delay_alu instid0(VALU_DEP_4) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_cndmask_b32_e32 v6, -1, v6, vcc_lo
	v_sub_co_u32 v8, vcc_lo, v3, 35
	v_subrev_co_ci_u32_e32 v9, vcc_lo, 0, v4, vcc_lo
	v_cmp_ne_u32_e32 vcc_lo, 0, v6
	v_cndmask_b32_e64 v6, -1, v7, s0
	s_add_u32 s0, s2, s4
	s_addc_u32 s1, s3, s5
	v_dual_cndmask_b32 v3, v3, v8 :: v_dual_cndmask_b32 v4, v4, v9
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_cmp_ne_u32_e32 vcc_lo, 0, v6
	v_dual_cndmask_b32 v1, v1, v3 :: v_dual_cndmask_b32 v2, v2, v4
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_xor_b32_e32 v1, v1, v5
	v_xor_b32_e32 v2, v2, v5
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_sub_co_u32 v1, vcc_lo, v1, v5
	v_sub_co_ci_u32_e32 v2, vcc_lo, v2, v5, vcc_lo
	global_store_b64 v0, v[1:2], s[0:1]
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z17entschluessselungPlS_
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 16
		.amdhsa_user_sgpr_count 15
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 10
		.amdhsa_next_free_sgpr 16
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end1:
	.size	_Z17entschluessselungPlS_, .Lfunc_end1-_Z17entschluessselungPlS_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 572
; NumSgprs: 18
; NumVgprs: 10
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 18
; NumVGPRsForWavesPerEU: 10
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 16
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z17verschluessselungPlS_
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z17verschluessselungPlS_.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     10
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 16
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z17entschluessselungPlS_
    .private_segment_fixed_size: 0
    .sgpr_count:     18
    .sgpr_spill_count: 0
    .symbol:         _Z17entschluessselungPlS_.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     10
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
