/*
 *  RFC 1321 compliant MD5 implementation
 *
 *  Copyright (C) 2006-2007  Christophe Devine
 *
 *  This library is free software; you can redistribute it and/or
 *  modify it under the terms of the GNU Lesser General Public
 *  License, version 2.1 as published by the Free Software Foundation.
 *
 *  This library is distributed in the hope that it will be useful,
 *  but WITHOUT ANY WARRANTY; without even the implied warranty of
 *  MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
 *  Lesser General Public License for more details.
 *
 *  You should have received a copy of the GNU Lesser General Public
 *  License along with this library; if not, write to the Free Software
 *  Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston,
 *  MA  02110-1301  USA
 */
/*
 *  The MD5 algorithm was designed by Ron Rivest in 1991.
 *
 *  http://www.ietf.org/rfc/rfc1321.txt
 */

#ifndef _CRT_SECURE_NO_DEPRECATE
#define _CRT_SECURE_NO_DEPRECATE 1
#endif

#include <string.h>
#include <stdio.h>

#include "md5.cuh"

/*
 * 32-bit integer manipulation macros (little endian)
 */
#ifndef GET_UINT32_LE
#define GET_UINT32_LE(n,b,i)                            \
{                                                       \
    (n) = ( (unsigned long) (b)[(i)    ]       )        \
        | ( (unsigned long) (b)[(i) + 1] <<  8 )        \
        | ( (unsigned long) (b)[(i) + 2] << 16 )        \
        | ( (unsigned long) (b)[(i) + 3] << 24 );       \
}
#endif

#ifndef PUT_UINT32_LE
#define PUT_UINT32_LE(n,b,i)                            \
{                                                       \
    (b)[(i)    ] = (unsigned char) ( (n)       );       \
    (b)[(i) + 1] = (unsigned char) ( (n) >>  8 );       \
    (b)[(i) + 2] = (unsigned char) ( (n) >> 16 );       \
    (b)[(i) + 3] = (unsigned char) ( (n) >> 24 );       \
}
#endif

/*
 * MD5 context setup
 */
__device__
void md5_starts(md5_context *ctx) {
	ctx->total[0] = 0;
	ctx->total[1] = 0;

	ctx->state[0] = 0x67452301;
	ctx->state[1] = 0xEFCDAB89;
	ctx->state[2] = 0x98BADCFE;
	ctx->state[3] = 0x10325476;
}

__device__
static void md5_process(md5_context *ctx, unsigned char data[64]) {
	unsigned long X[16], A, B, C, D;

	GET_UINT32_LE( X[ 0], data, 0 );
	GET_UINT32_LE( X[ 1], data, 4 );
	GET_UINT32_LE( X[ 2], data, 8 );
	GET_UINT32_LE( X[ 3], data, 12 );
	GET_UINT32_LE( X[ 4], data, 16 );
	GET_UINT32_LE( X[ 5], data, 20 );
	GET_UINT32_LE( X[ 6], data, 24 );
	GET_UINT32_LE( X[ 7], data, 28 );
	GET_UINT32_LE( X[ 8], data, 32 );
	GET_UINT32_LE( X[ 9], data, 36 );
	GET_UINT32_LE( X[10], data, 40 );
	GET_UINT32_LE( X[11], data, 44 );
	GET_UINT32_LE( X[12], data, 48 );
	GET_UINT32_LE( X[13], data, 52 );
	GET_UINT32_LE( X[14], data, 56 );
	GET_UINT32_LE( X[15], data, 60 );

#define S(x,n) ((x << n) | ((x & 0xFFFFFFFF) >> (32 - n)))

#define P(a,b,c,d,k,s,t)                                \
{                                                       \
    a += F(b,c,d) + X[k] + t; a = S(a,s) + b;           \
}

	A = ctx->state[0];
	B = ctx->state[1];
	C = ctx->state[2];
	D = ctx->state[3];

#define F(x,y,z) (z ^ (x & (y ^ z)))

	P( A, B, C, D, 0, 7, 0xD76AA478 );
	P( D, A, B, C, 1, 12, 0xE8C7B756 );
	P( C, D, A, B, 2, 17, 0x242070DB );
	P( B, C, D, A, 3, 22, 0xC1BDCEEE );
	P( A, B, C, D, 4, 7, 0xF57C0FAF );
	P( D, A, B, C, 5, 12, 0x4787C62A );
	P( C, D, A, B, 6, 17, 0xA8304613 );
	P( B, C, D, A, 7, 22, 0xFD469501 );
	P( A, B, C, D, 8, 7, 0x698098D8 );
	P( D, A, B, C, 9, 12, 0x8B44F7AF );
	P( C, D, A, B, 10, 17, 0xFFFF5BB1 );
	P( B, C, D, A, 11, 22, 0x895CD7BE );
	P( A, B, C, D, 12, 7, 0x6B901122 );
	P( D, A, B, C, 13, 12, 0xFD987193 );
	P( C, D, A, B, 14, 17, 0xA679438E );
	P( B, C, D, A, 15, 22, 0x49B40821 );

#undef F

#define F(x,y,z) (y ^ (z & (x ^ y)))

	P( A, B, C, D, 1, 5, 0xF61E2562 );
	P( D, A, B, C, 6, 9, 0xC040B340 );
	P( C, D, A, B, 11, 14, 0x265E5A51 );
	P( B, C, D, A, 0, 20, 0xE9B6C7AA );
	P( A, B, C, D, 5, 5, 0xD62F105D );
	P( D, A, B, C, 10, 9, 0x02441453 );
	P( C, D, A, B, 15, 14, 0xD8A1E681 );
	P( B, C, D, A, 4, 20, 0xE7D3FBC8 );
	P( A, B, C, D, 9, 5, 0x21E1CDE6 );
	P( D, A, B, C, 14, 9, 0xC33707D6 );
	P( C, D, A, B, 3, 14, 0xF4D50D87 );
	P( B, C, D, A, 8, 20, 0x455A14ED );
	P( A, B, C, D, 13, 5, 0xA9E3E905 );
	P( D, A, B, C, 2, 9, 0xFCEFA3F8 );
	P( C, D, A, B, 7, 14, 0x676F02D9 );
	P( B, C, D, A, 12, 20, 0x8D2A4C8A );

#undef F

#define F(x,y,z) (x ^ y ^ z)

	P( A, B, C, D, 5, 4, 0xFFFA3942 );
	P( D, A, B, C, 8, 11, 0x8771F681 );
	P( C, D, A, B, 11, 16, 0x6D9D6122 );
	P( B, C, D, A, 14, 23, 0xFDE5380C );
	P( A, B, C, D, 1, 4, 0xA4BEEA44 );
	P( D, A, B, C, 4, 11, 0x4BDECFA9 );
	P( C, D, A, B, 7, 16, 0xF6BB4B60 );
	P( B, C, D, A, 10, 23, 0xBEBFBC70 );
	P( A, B, C, D, 13, 4, 0x289B7EC6 );
	P( D, A, B, C, 0, 11, 0xEAA127FA );
	P( C, D, A, B, 3, 16, 0xD4EF3085 );
	P( B, C, D, A, 6, 23, 0x04881D05 );
	P( A, B, C, D, 9, 4, 0xD9D4D039 );
	P( D, A, B, C, 12, 11, 0xE6DB99E5 );
	P( C, D, A, B, 15, 16, 0x1FA27CF8 );
	P( B, C, D, A, 2, 23, 0xC4AC5665 );

#undef F

#define F(x,y,z) (y ^ (x | ~z))

	P( A, B, C, D, 0, 6, 0xF4292244 );
	P( D, A, B, C, 7, 10, 0x432AFF97 );
	P( C, D, A, B, 14, 15, 0xAB9423A7 );
	P( B, C, D, A, 5, 21, 0xFC93A039 );
	P( A, B, C, D, 12, 6, 0x655B59C3 );
	P( D, A, B, C, 3, 10, 0x8F0CCC92 );
	P( C, D, A, B, 10, 15, 0xFFEFF47D );
	P( B, C, D, A, 1, 21, 0x85845DD1 );
	P( A, B, C, D, 8, 6, 0x6FA87E4F );
	P( D, A, B, C, 15, 10, 0xFE2CE6E0 );
	P( C, D, A, B, 6, 15, 0xA3014314 );
	P( B, C, D, A, 13, 21, 0x4E0811A1 );
	P( A, B, C, D, 4, 6, 0xF7537E82 );
	P( D, A, B, C, 11, 10, 0xBD3AF235 );
	P( C, D, A, B, 2, 15, 0x2AD7D2BB );
	P( B, C, D, A, 9, 21, 0xEB86D391 );

#undef F
#undef P

	ctx->state[0] += A;
	ctx->state[1] += B;
	ctx->state[2] += C;
	ctx->state[3] += D;
}

/*
 * MD5 process buffer
 */
__device__
void md5_update(md5_context *ctx, unsigned char *input, int ilen) {
	int fill;
	unsigned long left;

	if (ilen <= 0)
		return;

	left = ctx->total[0] & 0x3F;
	fill = 64 - left;

	ctx->total[0] += ilen;
	ctx->total[0] &= 0xFFFFFFFF;

	if (ctx->total[0] < (unsigned long) ilen)
		ctx->total[1]++;

	if (left && ilen >= fill) {
		//memcpy((void *) (ctx->buffer + left), (void *) input, fill);
		unsigned char* toFill = ctx->buffer + left;
		for (int i = 0; i < fill; i++) {
			toFill[i] = input[i];
		}
		md5_process(ctx, ctx->buffer);
		input += fill;
		ilen -= fill;
		left = 0;
	}

	while (ilen >= 64) {
		md5_process(ctx, input);
		input += 64;
		ilen -= 64;
	}

	if (ilen > 0) {
		//memcpy((void *) (ctx->buffer + left), (void *) input, ilen);
		unsigned char* toFill = ctx->buffer + left;
		for (int i = 0; i < ilen; i++) {
			toFill[i] = input[i];
		}
	}
}

__device__
static const unsigned char md5_padding[64] = { 0x80, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
		0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
		0, 0, 0, 0 };

/*
 * MD5 final digest
 */
__device__
void md5_finish(md5_context *ctx, unsigned char *output) {
	unsigned long last, padn;
	unsigned long high, low;
	unsigned char msglen[8];

	high = (ctx->total[0] >> 29) | (ctx->total[1] << 3);
	low = (ctx->total[0] << 3);

	PUT_UINT32_LE( low, msglen, 0 );
	PUT_UINT32_LE( high, msglen, 4 );

	last = ctx->total[0] & 0x3F;
	padn = (last < 56) ? (56 - last) : (120 - last);

	md5_update(ctx, (unsigned char *) md5_padding, padn);
	md5_update(ctx, msglen, 8);

	PUT_UINT32_LE( ctx->state[0], output, 0 );
#ifndef FEATURE_REDUCED_HASH_SIZE
	PUT_UINT32_LE( ctx->state[1], output, 4 );
	PUT_UINT32_LE( ctx->state[2], output, 8 );
	PUT_UINT32_LE( ctx->state[3], output, 12 );
#endif
}

/*
 * Output = MD5( input buffer )
 */
__device__
void md5(unsigned char *input, int ilen, unsigned char *output) {
	md5_context ctx;

	md5_starts(&ctx);
	md5_update(&ctx, input, ilen);
	md5_finish(&ctx, output);

	//TODO What for?
	//memset(&ctx, 0, sizeof(md5_context));
}

	.file	"tmpxft_002b5fb7_00000000-6_md5.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2034:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2034:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z10md5_startsP11md5_context
	.type	_Z10md5_startsP11md5_context, @function
_Z10md5_startsP11md5_context:
.LFB2027:
	.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
.LFE2027:
	.size	_Z10md5_startsP11md5_context, .-_Z10md5_startsP11md5_context
	.globl	_Z10md5_updateP11md5_contextPhi
	.type	_Z10md5_updateP11md5_contextPhi, @function
_Z10md5_updateP11md5_contextPhi:
.LFB2029:
	.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
.LFE2029:
	.size	_Z10md5_updateP11md5_contextPhi, .-_Z10md5_updateP11md5_contextPhi
	.globl	_Z10md5_finishP11md5_contextPh
	.type	_Z10md5_finishP11md5_contextPh, @function
_Z10md5_finishP11md5_contextPh:
.LFB2030:
	.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
.LFE2030:
	.size	_Z10md5_finishP11md5_contextPh, .-_Z10md5_finishP11md5_contextPh
	.globl	_Z3md5PhiS_
	.type	_Z3md5PhiS_, @function
_Z3md5PhiS_:
.LFB2031:
	.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
.LFE2031:
	.size	_Z3md5PhiS_, .-_Z3md5PhiS_
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"md5_padding"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2057:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	pushq	$0
	.cfi_def_cfa_offset 24
	movl	$64, %r9d
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC0(%rip), %rdx
	movq	%rax, %rdi
	leaq	_ZL11md5_padding(%rip), %rsi
	movq	%rdx, %rcx
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterVar@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	popq	%rdx
	.cfi_def_cfa_offset 24
	popq	%rcx
	.cfi_def_cfa_offset 16
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rsi
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2057:
	.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 32
	.type	_ZL11md5_padding, @object
	.size	_ZL11md5_padding, 64
_ZL11md5_padding:
	.string	"\200"
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.string	""
	.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:

/*
 *  RFC 1321 compliant MD5 implementation
 *
 *  Copyright (C) 2006-2007  Christophe Devine
 *
 *  This library is free software; you can redistribute it and/or
 *  modify it under the terms of the GNU Lesser General Public
 *  License, version 2.1 as published by the Free Software Foundation.
 *
 *  This library is distributed in the hope that it will be useful,
 *  but WITHOUT ANY WARRANTY; without even the implied warranty of
 *  MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
 *  Lesser General Public License for more details.
 *
 *  You should have received a copy of the GNU Lesser General Public
 *  License along with this library; if not, write to the Free Software
 *  Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston,
 *  MA  02110-1301  USA
 */
/*
 *  The MD5 algorithm was designed by Ron Rivest in 1991.
 *
 *  http://www.ietf.org/rfc/rfc1321.txt
 */

#ifndef _CRT_SECURE_NO_DEPRECATE
#define _CRT_SECURE_NO_DEPRECATE 1
#endif


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

#include "md5.cuh"

/*
 * 32-bit integer manipulation macros (little endian)
 */
#ifndef GET_UINT32_LE
#define GET_UINT32_LE(n,b,i)                            \
{                                                       \
    (n) = ( (unsigned long) (b)[(i)    ]       )        \
        | ( (unsigned long) (b)[(i) + 1] <<  8 )        \
        | ( (unsigned long) (b)[(i) + 2] << 16 )        \
        | ( (unsigned long) (b)[(i) + 3] << 24 );       \
}
#endif

#ifndef PUT_UINT32_LE
#define PUT_UINT32_LE(n,b,i)                            \
{                                                       \
    (b)[(i)    ] = (unsigned char) ( (n)       );       \
    (b)[(i) + 1] = (unsigned char) ( (n) >>  8 );       \
    (b)[(i) + 2] = (unsigned char) ( (n) >> 16 );       \
    (b)[(i) + 3] = (unsigned char) ( (n) >> 24 );       \
}
#endif

/*
 * MD5 context setup
 */
__device__
void md5_starts(md5_context *ctx) {
	ctx->total[0] = 0;
	ctx->total[1] = 0;

	ctx->state[0] = 0x67452301;
	ctx->state[1] = 0xEFCDAB89;
	ctx->state[2] = 0x98BADCFE;
	ctx->state[3] = 0x10325476;
}

__device__
static void md5_process(md5_context *ctx, unsigned char data[64]) {
	unsigned long X[16], A, B, C, D;

	GET_UINT32_LE( X[ 0], data, 0 );
	GET_UINT32_LE( X[ 1], data, 4 );
	GET_UINT32_LE( X[ 2], data, 8 );
	GET_UINT32_LE( X[ 3], data, 12 );
	GET_UINT32_LE( X[ 4], data, 16 );
	GET_UINT32_LE( X[ 5], data, 20 );
	GET_UINT32_LE( X[ 6], data, 24 );
	GET_UINT32_LE( X[ 7], data, 28 );
	GET_UINT32_LE( X[ 8], data, 32 );
	GET_UINT32_LE( X[ 9], data, 36 );
	GET_UINT32_LE( X[10], data, 40 );
	GET_UINT32_LE( X[11], data, 44 );
	GET_UINT32_LE( X[12], data, 48 );
	GET_UINT32_LE( X[13], data, 52 );
	GET_UINT32_LE( X[14], data, 56 );
	GET_UINT32_LE( X[15], data, 60 );

#define S(x,n) ((x << n) | ((x & 0xFFFFFFFF) >> (32 - n)))

#define P(a,b,c,d,k,s,t)                                \
{                                                       \
    a += F(b,c,d) + X[k] + t; a = S(a,s) + b;           \
}

	A = ctx->state[0];
	B = ctx->state[1];
	C = ctx->state[2];
	D = ctx->state[3];

#define F(x,y,z) (z ^ (x & (y ^ z)))

	P( A, B, C, D, 0, 7, 0xD76AA478 );
	P( D, A, B, C, 1, 12, 0xE8C7B756 );
	P( C, D, A, B, 2, 17, 0x242070DB );
	P( B, C, D, A, 3, 22, 0xC1BDCEEE );
	P( A, B, C, D, 4, 7, 0xF57C0FAF );
	P( D, A, B, C, 5, 12, 0x4787C62A );
	P( C, D, A, B, 6, 17, 0xA8304613 );
	P( B, C, D, A, 7, 22, 0xFD469501 );
	P( A, B, C, D, 8, 7, 0x698098D8 );
	P( D, A, B, C, 9, 12, 0x8B44F7AF );
	P( C, D, A, B, 10, 17, 0xFFFF5BB1 );
	P( B, C, D, A, 11, 22, 0x895CD7BE );
	P( A, B, C, D, 12, 7, 0x6B901122 );
	P( D, A, B, C, 13, 12, 0xFD987193 );
	P( C, D, A, B, 14, 17, 0xA679438E );
	P( B, C, D, A, 15, 22, 0x49B40821 );

#undef F

#define F(x,y,z) (y ^ (z & (x ^ y)))

	P( A, B, C, D, 1, 5, 0xF61E2562 );
	P( D, A, B, C, 6, 9, 0xC040B340 );
	P( C, D, A, B, 11, 14, 0x265E5A51 );
	P( B, C, D, A, 0, 20, 0xE9B6C7AA );
	P( A, B, C, D, 5, 5, 0xD62F105D );
	P( D, A, B, C, 10, 9, 0x02441453 );
	P( C, D, A, B, 15, 14, 0xD8A1E681 );
	P( B, C, D, A, 4, 20, 0xE7D3FBC8 );
	P( A, B, C, D, 9, 5, 0x21E1CDE6 );
	P( D, A, B, C, 14, 9, 0xC33707D6 );
	P( C, D, A, B, 3, 14, 0xF4D50D87 );
	P( B, C, D, A, 8, 20, 0x455A14ED );
	P( A, B, C, D, 13, 5, 0xA9E3E905 );
	P( D, A, B, C, 2, 9, 0xFCEFA3F8 );
	P( C, D, A, B, 7, 14, 0x676F02D9 );
	P( B, C, D, A, 12, 20, 0x8D2A4C8A );

#undef F

#define F(x,y,z) (x ^ y ^ z)

	P( A, B, C, D, 5, 4, 0xFFFA3942 );
	P( D, A, B, C, 8, 11, 0x8771F681 );
	P( C, D, A, B, 11, 16, 0x6D9D6122 );
	P( B, C, D, A, 14, 23, 0xFDE5380C );
	P( A, B, C, D, 1, 4, 0xA4BEEA44 );
	P( D, A, B, C, 4, 11, 0x4BDECFA9 );
	P( C, D, A, B, 7, 16, 0xF6BB4B60 );
	P( B, C, D, A, 10, 23, 0xBEBFBC70 );
	P( A, B, C, D, 13, 4, 0x289B7EC6 );
	P( D, A, B, C, 0, 11, 0xEAA127FA );
	P( C, D, A, B, 3, 16, 0xD4EF3085 );
	P( B, C, D, A, 6, 23, 0x04881D05 );
	P( A, B, C, D, 9, 4, 0xD9D4D039 );
	P( D, A, B, C, 12, 11, 0xE6DB99E5 );
	P( C, D, A, B, 15, 16, 0x1FA27CF8 );
	P( B, C, D, A, 2, 23, 0xC4AC5665 );

#undef F

#define F(x,y,z) (y ^ (x | ~z))

	P( A, B, C, D, 0, 6, 0xF4292244 );
	P( D, A, B, C, 7, 10, 0x432AFF97 );
	P( C, D, A, B, 14, 15, 0xAB9423A7 );
	P( B, C, D, A, 5, 21, 0xFC93A039 );
	P( A, B, C, D, 12, 6, 0x655B59C3 );
	P( D, A, B, C, 3, 10, 0x8F0CCC92 );
	P( C, D, A, B, 10, 15, 0xFFEFF47D );
	P( B, C, D, A, 1, 21, 0x85845DD1 );
	P( A, B, C, D, 8, 6, 0x6FA87E4F );
	P( D, A, B, C, 15, 10, 0xFE2CE6E0 );
	P( C, D, A, B, 6, 15, 0xA3014314 );
	P( B, C, D, A, 13, 21, 0x4E0811A1 );
	P( A, B, C, D, 4, 6, 0xF7537E82 );
	P( D, A, B, C, 11, 10, 0xBD3AF235 );
	P( C, D, A, B, 2, 15, 0x2AD7D2BB );
	P( B, C, D, A, 9, 21, 0xEB86D391 );

#undef F
#undef P

	ctx->state[0] += A;
	ctx->state[1] += B;
	ctx->state[2] += C;
	ctx->state[3] += D;
}

/*
 * MD5 process buffer
 */
__device__
void md5_update(md5_context *ctx, unsigned char *input, int ilen) {
	int fill;
	unsigned long left;

	if (ilen <= 0)
		return;

	left = ctx->total[0] & 0x3F;
	fill = 64 - left;

	ctx->total[0] += ilen;
	ctx->total[0] &= 0xFFFFFFFF;

	if (ctx->total[0] < (unsigned long) ilen)
		ctx->total[1]++;

	if (left && ilen >= fill) {
		//memcpy((void *) (ctx->buffer + left), (void *) input, fill);
		unsigned char* toFill = ctx->buffer + left;
		for (int i = 0; i < fill; i++) {
			toFill[i] = input[i];
		}
		md5_process(ctx, ctx->buffer);
		input += fill;
		ilen -= fill;
		left = 0;
	}

	while (ilen >= 64) {
		md5_process(ctx, input);
		input += 64;
		ilen -= 64;
	}

	if (ilen > 0) {
		//memcpy((void *) (ctx->buffer + left), (void *) input, ilen);
		unsigned char* toFill = ctx->buffer + left;
		for (int i = 0; i < ilen; i++) {
			toFill[i] = input[i];
		}
	}
}

__device__
static const unsigned char md5_padding[64] = { 0x80, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
		0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
		0, 0, 0, 0 };

/*
 * MD5 final digest
 */
__device__
void md5_finish(md5_context *ctx, unsigned char *output) {
	unsigned long last, padn;
	unsigned long high, low;
	unsigned char msglen[8];

	high = (ctx->total[0] >> 29) | (ctx->total[1] << 3);
	low = (ctx->total[0] << 3);

	PUT_UINT32_LE( low, msglen, 0 );
	PUT_UINT32_LE( high, msglen, 4 );

	last = ctx->total[0] & 0x3F;
	padn = (last < 56) ? (56 - last) : (120 - last);

	md5_update(ctx, (unsigned char *) md5_padding, padn);
	md5_update(ctx, msglen, 8);

	PUT_UINT32_LE( ctx->state[0], output, 0 );
#ifndef FEATURE_REDUCED_HASH_SIZE
	PUT_UINT32_LE( ctx->state[1], output, 4 );
	PUT_UINT32_LE( ctx->state[2], output, 8 );
	PUT_UINT32_LE( ctx->state[3], output, 12 );
#endif
}

/*
 * Output = MD5( input buffer )
 */
__device__
void md5(unsigned char *input, int ilen, unsigned char *output) {
	md5_context ctx;

	md5_starts(&ctx);
	md5_update(&ctx, input, ilen);
	md5_finish(&ctx, output);

	//TODO What for?
	//memset(&ctx, 0, sizeof(md5_context));
}

	.text
	.file	"md5.hip"
	.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 __hip_cuid_

	.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_
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.amdgpu_metadata
---
amdhsa.kernels:  []
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
