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

#include <cuda_runtime_api.h>
#include <cuda.h>

#define START_RED printf("\033[1;31m");
#define START_GRN printf("\033[1;32m");
#define START_YLW printf("\033[1;33m");
#define START_BLU printf("\033[1;34m");
#define END printf("\033[0m");

#define ONE_SEC (1000 * 1000 * 1000)
#define HALF (1024 * 32 * 64)

#define NORM 0

#define LOOP 0

#define SEL 1

#define CASE 17

#define REVERSE 0

FILE * data = fopen("data.txt", "r");

uint64_t latency[17] = { 0 };
const char* size_str[17] = { "64", "128", "256", "512", "1024", "1514",\
		 "64 * 32", "64 * 64", "64 * 128", "64 * 256",	\
		"64 * 512", "64 * 1024", "64 * 1024 * 2", "64 * 1024 * 4",		 	\
		"64 * 1024 * 8", "64 * 1024 * 16", "64 * 1024 * 32"};
int size[17] = { 64, 128, 256, 512, 1024, 1514, 64 * 32, 64 * 64,\
			 64 * 128, 64 * 256, 64 * 512, 64 * 1024,\
			 64 * 1024 * 2, 64 * 1024 * 4, 64 * 1024 * 8,\
			 64 * 1024 * 16, 64 * 1024 * 32};

int start[17] = { 0 };
int end[17] = { 0 };

int monotonic_time() 
{
	struct timespec timespec;
	clock_gettime(CLOCK_MONOTONIC, &timespec);
	return timespec.tv_sec * ONE_SEC + timespec.tv_nsec;
}

void call_data(char * host_buf, int size)
{
	fseek(data, 0, SEEK_SET);
	fgets(host_buf, sizeof(char) * size, data);
/*
	for(int i = 0; i < size; i++)
	{
		fscanf(data, "%c", host_buf + i);
	}
*/
}

void once(char * device_buf, char * host_buf, int test_cnt)
{
	int i = 0;
	
	int skip = 0;

	while(i < test_cnt)
	{
		call_data(host_buf, HALF * 2);
		skip = 0;
#if REVERSE
		for(int j = CASE - 1; j >= 0; j--)
#else
		for(int j = 0; j < CASE; j++)
#endif
		{
#if NORM
			skip += size[j];
#else
			skip = rand();
#endif
			if(skip >= HALF)
				skip %= HALF;
			start[j] = monotonic_time();
			cudaMemcpy(device_buf, host_buf + skip, size[j], cudaMemcpyHostToDevice);
			end[j] = monotonic_time();
			latency[j] += end[j] - start[j];
			cudaMemset(device_buf, 0, size[j]);
		}
		i++;
	}

	for(i = 0; i < CASE; i++)
	{
		latency[i] /= (uint64_t)test_cnt;
	}
}

void loop(char * device_buf, char * host_buf, int test_cnt, int loop_cnt)
{
	int i = 0;

	int skip = 0;

	while(i < test_cnt)
	{
		call_data(host_buf, HALF * 2);
		skip = 0;
#if REVERSE
		for(int j = CASE - 1; j >= 0; j--)
#else
		for(int j = 0; j < CASE; j++)
#endif
		{
#if NORM
			skip += size[j];
#else
			skip = rand();
#endif
			if(skip >= HALF)
				skip %= HALF;
			start[j] = monotonic_time();
			for(int k = 0; k < loop_cnt; k++)
			{
			cudaMemcpy(device_buf, host_buf + skip, size[j], cudaMemcpyHostToDevice);
			}
			end[j] = monotonic_time();
			latency[j] += (end[j] - start[j]) / (uint64_t)loop_cnt;
		}
		i++;
	}

	for(i = 0; i < CASE; i++)
	{
		latency[i] /= (uint64_t)test_cnt;
	}
}

void same_cnt_loop(char * device_buf, char * host_buf, int test_cnt)
{
	int i = 0;

	int loop_cnt = size[16];
	int cur_loop_cnt;
	int skip = 0;

	while(i < test_cnt)
	{
		call_data(host_buf, HALF * 2);
		skip = 0;
#if REVERSE
		for(int j = CASE - 1; j >= 0; j--)
#else
		for(int j = 0; j < CASE; j++)
#endif
		{
			cur_loop_cnt = loop_cnt / size[j];
#if NORM
			skip += size[j];
#else
			skip = rand();
#endif
			if(skip >= HALF)
				skip %= HALF;
			start[j] = monotonic_time();
			for(int k = 0; k < cur_loop_cnt; k++)
			{
			cudaMemcpy(device_buf, host_buf + skip, size[j], cudaMemcpyHostToDevice);
			}
			end[j] = monotonic_time();
			latency[j] += end[j] - start[j];
		}
		i++;
	}

	for(i = 0; i < CASE; i++)
	{
		latency[i] /= (uint64_t)test_cnt;
	}
}

void print_result(int test_cnt)
{
	START_RED
	printf("\n\n______________________TEST START_______________________\n\n");
	END

	START_YLW
#if NORM
	printf("   NORMAL DATA TEST!\n");
#else
	printf("   RANDOM DATA TEST!\n");
#endif
	END

	START_GRN
#if LOOP
	printf("   TEST WAS RUNNED %d TIMES!\n", test_cnt);
#elif SEL
	printf("   TEST WAS RUNNED SAME TIMES!\n");
#else
	printf("   TEST WAS RUNNED ONCE!\n");
#endif
	END

#if REVERSE
	START_BLU
	printf("   TEST WAS RUNNED REVERSED ORDER!\n");
	END
#endif

	for(int i = 0; i < CASE; i++)
	{
		printf("   data size : %s, latency : %ld\n", size_str[i], latency[i]);
	}

	START_RED
	printf("\n______________________TEST END_______________________\n\n\n");
	END
}


int main(void)
{
	int test_cnt;

	char * device_buf;
	char * host_buf;

	srand(time(NULL));

	host_buf = (char *)calloc(HALF * 2, sizeof(char));
	cudaHostAlloc((void**)&device_buf, HALF * sizeof(char), cudaHostAllocDefault);
	cudaMemset(device_buf, 0 ,HALF * sizeof(char)); 

	printf("Enter the test_cnt\n");
	scanf("%d", &test_cnt);

#if LOOP
	int loop_cnt;

	printf("Enter the loop_cnt\n");
	scanf("%d", &loop_cnt);
	loop(device_buf, host_buf, test_cnt, loop_cnt);
#elif SEL
	same_cnt_loop(device_buf, host_buf, test_cnt);
#else
	once(device_buf, host_buf, test_cnt);
#endif
	print_result(test_cnt);

	cudaFree(device_buf);

	fclose(data);

	return 0;
}

	.file	"tmpxft_002ea2d1_00000000-6_memcpy_test.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2036:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2036:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z14monotonic_timev
	.type	_Z14monotonic_timev, @function
_Z14monotonic_timev:
.LFB2027:
	.cfi_startproc
	endbr64
	subq	$40, %rsp
	.cfi_def_cfa_offset 48
	movl	$1, %edi
	movq	%fs:40, %rax
	movq	%rax, 24(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rsi
	call	clock_gettime@PLT
	imull	$1000000000, 8(%rsp), %eax
	addl	16(%rsp), %eax
	movq	24(%rsp), %rdx
	subq	%fs:40, %rdx
	je	.L3
	call	__stack_chk_fail@PLT
.L3:
	addq	$40, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2027:
	.size	_Z14monotonic_timev, .-_Z14monotonic_timev
	.globl	_Z9call_dataPci
	.type	_Z9call_dataPci, @function
_Z9call_dataPci:
.LFB2028:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	movl	%esi, %ebp
	xorl	%edx, %edx
	xorl	%esi, %esi
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	movq	%rdi, %rbx
	pushq	%rax
	.cfi_def_cfa_offset 32
	movq	data(%rip), %rdi
	call	fseek@PLT
	movl	%ebp, %edx
	movq	%rbx, %rdi
	orq	$-1, %rsi
	movq	data(%rip), %rcx
	popq	%r8
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	jmp	__fgets_chk@PLT
	.cfi_endproc
.LFE2028:
	.size	_Z9call_dataPci, .-_Z9call_dataPci
	.globl	_Z4oncePcS_i
	.type	_Z4oncePcS_i, @function
_Z4oncePcS_i:
.LFB2029:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	xorl	%r14d, %r14d
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movslq	%edx, %rbx
	subq	$40, %rsp
	.cfi_def_cfa_offset 96
	movq	%rdi, 8(%rsp)
	movq	%rsi, 16(%rsp)
.L9:
	cmpl	%ebx, %r14d
	jge	.L17
	movq	16(%rsp), %rdi
	movl	$4194304, %esi
	leaq	start(%rip), %rbp
	leaq	size(%rip), %r13
	leaq	end(%rip), %r15
	leaq	latency(%rip), %r12
	call	_Z9call_dataPci
.L11:
	call	rand@PLT
	movl	%eax, %esi
	cmpl	$2097151, %eax
	jle	.L10
	andl	$2097151, %esi
.L10:
	movl	%esi, 28(%rsp)
	addq	$4, %rbp
	addq	$4, %r13
	addq	$4, %r15
	call	_Z14monotonic_timev
	movslq	28(%rsp), %rsi
	movslq	-4(%r13), %rdx
	movl	$1, %ecx
	movl	%eax, -4(%rbp)
	movq	16(%rsp), %rax
	movq	8(%rsp), %rdi
	addq	%rax, %rsi
	call	cudaMemcpy@PLT
	call	_Z14monotonic_timev
	movslq	-4(%r13), %rdx
	movq	8(%rsp), %rdi
	xorl	%esi, %esi
	movl	%eax, -4(%r15)
	subl	-4(%rbp), %eax
	cltq
	addq	%rax, (%r12)
	addq	$8, %r12
	call	cudaMemset@PLT
	leaq	68+start(%rip), %rax
	cmpq	%rax, %rbp
	jne	.L11
	incl	%r14d
	jmp	.L9
.L17:
	leaq	latency(%rip), %rcx
	leaq	136(%rcx), %rsi
.L13:
	movq	(%rcx), %rax
	xorl	%edx, %edx
	addq	$8, %rcx
	divq	%rbx
	movq	%rax, -8(%rcx)
	cmpq	%rsi, %rcx
	jne	.L13
	addq	$40, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2029:
	.size	_Z4oncePcS_i, .-_Z4oncePcS_i
	.globl	_Z4loopPcS_ii
	.type	_Z4loopPcS_ii, @function
_Z4loopPcS_ii:
.LFB2030:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movslq	%ecx, %rax
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	xorl	%r13d, %r13d
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movq	%rax, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movq	%rsi, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movslq	%edx, %rbx
	subq	$40, %rsp
	.cfi_def_cfa_offset 96
	movq	%rdi, (%rsp)
	movq	%rax, 8(%rsp)
.L19:
	cmpl	%ebx, %r13d
	jge	.L29
	movl	$4194304, %esi
	movq	%rbp, %rdi
	xorl	%r14d, %r14d
	call	_Z9call_dataPci
.L23:
	call	rand@PLT
	movl	%eax, %r9d
	cmpl	$2097151, %eax
	jle	.L20
	andl	$2097151, %r9d
.L20:
	movl	%r9d, 16(%rsp)
	leaq	start(%rip), %r15
	call	_Z14monotonic_timev
	movslq	16(%rsp), %r9
	xorl	%r10d, %r10d
	movl	%eax, (%r15,%r14)
.L21:
	cmpl	%r12d, %r10d
	jge	.L30
	leaq	size(%rip), %rax
	movq	(%rsp), %rdi
	leaq	0(%rbp,%r9), %rsi
	movl	$1, %ecx
	movslq	(%rax,%r14), %rdx
	movl	%r10d, 28(%rsp)
	movq	%r9, 16(%rsp)
	call	cudaMemcpy@PLT
	movl	28(%rsp), %r10d
	movq	16(%rsp), %r9
	incl	%r10d
	jmp	.L21
.L30:
	call	_Z14monotonic_timev
	leaq	end(%rip), %rdx
	leaq	latency(%rip), %rcx
	movl	%eax, (%rdx,%r14)
	subl	(%r15,%r14), %eax
	xorl	%edx, %edx
	cltq
	divq	8(%rsp)
	addq	%rax, (%rcx,%r14,2)
	addq	$4, %r14
	cmpq	$68, %r14
	jne	.L23
	incl	%r13d
	jmp	.L19
.L29:
	leaq	latency(%rip), %rcx
	leaq	136(%rcx), %rsi
.L25:
	movq	(%rcx), %rax
	xorl	%edx, %edx
	addq	$8, %rcx
	divq	%rbx
	movq	%rax, -8(%rcx)
	cmpq	%rsi, %rcx
	jne	.L25
	addq	$40, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2030:
	.size	_Z4loopPcS_ii, .-_Z4loopPcS_ii
	.globl	_Z13same_cnt_loopPcS_i
	.type	_Z13same_cnt_loopPcS_i, @function
_Z13same_cnt_loopPcS_i:
.LFB2031:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	xorl	%r14d, %r14d
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movslq	%edx, %rbx
	subq	$56, %rsp
	.cfi_def_cfa_offset 112
	movl	64+size(%rip), %eax
	movq	%rdi, 24(%rsp)
	movq	%rsi, 8(%rsp)
	movl	%eax, 20(%rsp)
.L32:
	cmpl	%ebx, %r14d
	jge	.L42
	movq	8(%rsp), %rdi
	movl	$4194304, %esi
	leaq	size(%rip), %rbp
	leaq	end(%rip), %r15
	leaq	latency(%rip), %r13
	leaq	start(%rip), %r12
	call	_Z9call_dataPci
.L36:
	movl	20(%rsp), %eax
	cltd
	idivl	0(%rbp)
	movl	%eax, 40(%rsp)
	call	rand@PLT
	movl	%eax, %r8d
	cmpl	$2097151, %eax
	jle	.L33
	andl	$2097151, %r8d
.L33:
	movl	%r8d, 32(%rsp)
	call	_Z14monotonic_timev
	movslq	32(%rsp), %r8
	xorl	%r9d, %r9d
	movl	%eax, (%r12)
.L34:
	movl	40(%rsp), %eax
	cmpl	%eax, %r9d
	jge	.L43
	movq	8(%rsp), %rax
	movslq	0(%rbp), %rdx
	movl	$1, %ecx
	movl	%r9d, 44(%rsp)
	movq	24(%rsp), %rdi
	movq	%r8, 32(%rsp)
	leaq	(%rax,%r8), %rsi
	call	cudaMemcpy@PLT
	movl	44(%rsp), %r9d
	movq	32(%rsp), %r8
	incl	%r9d
	jmp	.L34
.L43:
	call	_Z14monotonic_timev
	addq	$4, %rbp
	addq	$4, %r15
	addq	$4, %r12
	movl	%eax, -4(%r15)
	subl	-4(%r12), %eax
	cltq
	addq	%rax, 0(%r13)
	leaq	68+size(%rip), %rax
	addq	$8, %r13
	cmpq	%rax, %rbp
	jne	.L36
	incl	%r14d
	jmp	.L32
.L42:
	leaq	latency(%rip), %rcx
	leaq	136(%rcx), %rsi
.L38:
	movq	(%rcx), %rax
	xorl	%edx, %edx
	addq	$8, %rcx
	divq	%rbx
	movq	%rax, -8(%rcx)
	cmpq	%rsi, %rcx
	jne	.L38
	addq	$56, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2031:
	.size	_Z13same_cnt_loopPcS_i, .-_Z13same_cnt_loopPcS_i
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"\033[1;31m"
.LC1:
	.string	"\n\n______________________TEST START_______________________\n\n"
.LC2:
	.string	"\033[0m"
.LC3:
	.string	"\033[1;33m"
.LC4:
	.string	"   RANDOM DATA TEST!\n"
.LC5:
	.string	"\033[1;32m"
.LC6:
	.string	"   TEST WAS RUNNED SAME TIMES!\n"
.LC7:
	.string	"   data size : %s, latency : %ld\n"
.LC8:
	.string	"\n______________________TEST END_______________________\n\n\n"
	.text
	.globl	_Z12print_resulti
	.type	_Z12print_resulti, @function
_Z12print_resulti:
.LFB2032:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movl	$2, %edi
	xorl	%eax, %eax
	leaq	.LC7(%rip), %r15
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	leaq	.LC0(%rip), %r13
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movq	%r13, %rsi
	leaq	size_str(%rip), %r12
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	leaq	latency(%rip), %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	leaq	.LC2(%rip), %rbx
	leaq	136(%rbp), %r14
	pushq	%rcx
	.cfi_def_cfa_offset 64
	call	__printf_chk@PLT
	leaq	.LC1(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	%rbx, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	leaq	.LC3(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	leaq	.LC4(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	%rbx, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	leaq	.LC5(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	leaq	.LC6(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	%rbx, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
.L45:
	movq	0(%rbp), %rcx
	movq	(%r12), %rdx
	movq	%r15, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	addq	$8, %rbp
	addq	$8, %r12
	call	__printf_chk@PLT
	cmpq	%r14, %rbp
	jne	.L45
	movq	%r13, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	leaq	.LC8(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	popq	%rdx
	.cfi_def_cfa_offset 56
	movq	%rbx, %rsi
	movl	$2, %edi
	popq	%rbx
	.cfi_def_cfa_offset 48
	xorl	%eax, %eax
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	jmp	__printf_chk@PLT
	.cfi_endproc
.LFE2032:
	.size	_Z12print_resulti, .-_Z12print_resulti
	.section	.rodata.str1.1
.LC9:
	.string	"Enter the test_cnt\n"
.LC10:
	.string	"%d"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2033:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	xorl	%edi, %edi
	subq	$32, %rsp
	.cfi_def_cfa_offset 48
	movq	%fs:40, %rax
	movq	%rax, 24(%rsp)
	xorl	%eax, %eax
	call	time@PLT
	movl	%eax, %edi
	call	srand@PLT
	movl	$1, %esi
	movl	$4194304, %edi
	call	calloc@PLT
	xorl	%edx, %edx
	leaq	16(%rsp), %rdi
	movl	$2097152, %esi
	movq	%rax, %rbx
	call	cudaHostAlloc@PLT
	movq	16(%rsp), %rdi
	movl	$2097152, %edx
	xorl	%esi, %esi
	call	cudaMemset@PLT
	leaq	.LC9(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	leaq	12(%rsp), %rsi
	leaq	.LC10(%rip), %rdi
	xorl	%eax, %eax
	call	__isoc23_scanf@PLT
	movl	12(%rsp), %edx
	movq	16(%rsp), %rdi
	movq	%rbx, %rsi
	call	_Z13same_cnt_loopPcS_i
	movl	12(%rsp), %edi
	call	_Z12print_resulti
	movq	16(%rsp), %rdi
	call	cudaFree@PLT
	movq	data(%rip), %rdi
	call	fclose@PLT
	movq	24(%rsp), %rax
	subq	%fs:40, %rax
	je	.L49
	call	__stack_chk_fail@PLT
.L49:
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	xorl	%eax, %eax
	popq	%rbx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2033:
	.size	main, .-main
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2059:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	movq	%rax, %rdi
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rdx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2059:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.rodata.str1.1
.LC11:
	.string	"r"
.LC12:
	.string	"data.txt"
	.section	.text.startup
	.type	_GLOBAL__sub_I_data, @function
_GLOBAL__sub_I_data:
.LFB2175:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	leaq	.LC11(%rip), %rsi
	leaq	.LC12(%rip), %rdi
	call	fopen@PLT
	movq	%rax, data(%rip)
	popq	%rdx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2175:
	.size	_GLOBAL__sub_I_data, .-_GLOBAL__sub_I_data
	.section	.init_array
	.align 8
	.quad	_GLOBAL__sub_I_data
	.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
	.globl	end
	.bss
	.align 32
	.type	end, @object
	.size	end, 68
end:
	.zero	68
	.globl	start
	.align 32
	.type	start, @object
	.size	start, 68
start:
	.zero	68
	.globl	size
	.data
	.align 32
	.type	size, @object
	.size	size, 68
size:
	.long	64
	.long	128
	.long	256
	.long	512
	.long	1024
	.long	1514
	.long	2048
	.long	4096
	.long	8192
	.long	16384
	.long	32768
	.long	65536
	.long	131072
	.long	262144
	.long	524288
	.long	1048576
	.long	2097152
	.globl	size_str
	.section	.rodata.str1.1
.LC13:
	.string	"64"
.LC14:
	.string	"128"
.LC15:
	.string	"256"
.LC16:
	.string	"512"
.LC17:
	.string	"1024"
.LC18:
	.string	"1514"
.LC19:
	.string	"64 * 32"
.LC20:
	.string	"64 * 64"
.LC21:
	.string	"64 * 128"
.LC22:
	.string	"64 * 256"
.LC23:
	.string	"64 * 512"
.LC24:
	.string	"64 * 1024"
.LC25:
	.string	"64 * 1024 * 2"
.LC26:
	.string	"64 * 1024 * 4"
.LC27:
	.string	"64 * 1024 * 8"
.LC28:
	.string	"64 * 1024 * 16"
.LC29:
	.string	"64 * 1024 * 32"
	.section	.data.rel.local,"aw"
	.align 32
	.type	size_str, @object
	.size	size_str, 136
size_str:
	.quad	.LC13
	.quad	.LC14
	.quad	.LC15
	.quad	.LC16
	.quad	.LC17
	.quad	.LC18
	.quad	.LC19
	.quad	.LC20
	.quad	.LC21
	.quad	.LC22
	.quad	.LC23
	.quad	.LC24
	.quad	.LC25
	.quad	.LC26
	.quad	.LC27
	.quad	.LC28
	.quad	.LC29
	.globl	latency
	.bss
	.align 32
	.type	latency, @object
	.size	latency, 136
latency:
	.zero	136
	.globl	data
	.align 8
	.type	data, @object
	.size	data, 8
data:
	.zero	8
	.ident	"GCC: (Ubuntu 13.3.0-6ubuntu2~24.04) 13.3.0"
	.section	.note.GNU-stack,"",@progbits
	.section	.note.gnu.property,"a"
	.align 8
	.long	1f - 0f
	.long	4f - 1f
	.long	5
0:
	.string	"GNU"
1:
	.align 8
	.long	0xc0000002
	.long	3f - 2f
2:
	.long	0x3
3:
	.align 8
4:

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

#include <hip/hip_runtime_api.h>
#include <hip/hip_runtime.h>

#define START_RED printf("\033[1;31m");
#define START_GRN printf("\033[1;32m");
#define START_YLW printf("\033[1;33m");
#define START_BLU printf("\033[1;34m");
#define END printf("\033[0m");

#define ONE_SEC (1000 * 1000 * 1000)
#define HALF (1024 * 32 * 64)

#define NORM 0

#define LOOP 0

#define SEL 1

#define CASE 17

#define REVERSE 0

FILE * data = fopen("data.txt", "r");

uint64_t latency[17] = { 0 };
const char* size_str[17] = { "64", "128", "256", "512", "1024", "1514",\
		 "64 * 32", "64 * 64", "64 * 128", "64 * 256",	\
		"64 * 512", "64 * 1024", "64 * 1024 * 2", "64 * 1024 * 4",		 	\
		"64 * 1024 * 8", "64 * 1024 * 16", "64 * 1024 * 32"};
int size[17] = { 64, 128, 256, 512, 1024, 1514, 64 * 32, 64 * 64,\
			 64 * 128, 64 * 256, 64 * 512, 64 * 1024,\
			 64 * 1024 * 2, 64 * 1024 * 4, 64 * 1024 * 8,\
			 64 * 1024 * 16, 64 * 1024 * 32};

int start[17] = { 0 };
int end[17] = { 0 };

int monotonic_time() 
{
	struct timespec timespec;
	clock_gettime(CLOCK_MONOTONIC, &timespec);
	return timespec.tv_sec * ONE_SEC + timespec.tv_nsec;
}

void call_data(char * host_buf, int size)
{
	fseek(data, 0, SEEK_SET);
	fgets(host_buf, sizeof(char) * size, data);
/*
	for(int i = 0; i < size; i++)
	{
		fscanf(data, "%c", host_buf + i);
	}
*/
}

void once(char * device_buf, char * host_buf, int test_cnt)
{
	int i = 0;
	
	int skip = 0;

	while(i < test_cnt)
	{
		call_data(host_buf, HALF * 2);
		skip = 0;
#if REVERSE
		for(int j = CASE - 1; j >= 0; j--)
#else
		for(int j = 0; j < CASE; j++)
#endif
		{
#if NORM
			skip += size[j];
#else
			skip = rand();
#endif
			if(skip >= HALF)
				skip %= HALF;
			start[j] = monotonic_time();
			hipMemcpy(device_buf, host_buf + skip, size[j], hipMemcpyHostToDevice);
			end[j] = monotonic_time();
			latency[j] += end[j] - start[j];
			hipMemset(device_buf, 0, size[j]);
		}
		i++;
	}

	for(i = 0; i < CASE; i++)
	{
		latency[i] /= (uint64_t)test_cnt;
	}
}

void loop(char * device_buf, char * host_buf, int test_cnt, int loop_cnt)
{
	int i = 0;

	int skip = 0;

	while(i < test_cnt)
	{
		call_data(host_buf, HALF * 2);
		skip = 0;
#if REVERSE
		for(int j = CASE - 1; j >= 0; j--)
#else
		for(int j = 0; j < CASE; j++)
#endif
		{
#if NORM
			skip += size[j];
#else
			skip = rand();
#endif
			if(skip >= HALF)
				skip %= HALF;
			start[j] = monotonic_time();
			for(int k = 0; k < loop_cnt; k++)
			{
			hipMemcpy(device_buf, host_buf + skip, size[j], hipMemcpyHostToDevice);
			}
			end[j] = monotonic_time();
			latency[j] += (end[j] - start[j]) / (uint64_t)loop_cnt;
		}
		i++;
	}

	for(i = 0; i < CASE; i++)
	{
		latency[i] /= (uint64_t)test_cnt;
	}
}

void same_cnt_loop(char * device_buf, char * host_buf, int test_cnt)
{
	int i = 0;

	int loop_cnt = size[16];
	int cur_loop_cnt;
	int skip = 0;

	while(i < test_cnt)
	{
		call_data(host_buf, HALF * 2);
		skip = 0;
#if REVERSE
		for(int j = CASE - 1; j >= 0; j--)
#else
		for(int j = 0; j < CASE; j++)
#endif
		{
			cur_loop_cnt = loop_cnt / size[j];
#if NORM
			skip += size[j];
#else
			skip = rand();
#endif
			if(skip >= HALF)
				skip %= HALF;
			start[j] = monotonic_time();
			for(int k = 0; k < cur_loop_cnt; k++)
			{
			hipMemcpy(device_buf, host_buf + skip, size[j], hipMemcpyHostToDevice);
			}
			end[j] = monotonic_time();
			latency[j] += end[j] - start[j];
		}
		i++;
	}

	for(i = 0; i < CASE; i++)
	{
		latency[i] /= (uint64_t)test_cnt;
	}
}

void print_result(int test_cnt)
{
	START_RED
	printf("\n\n______________________TEST START_______________________\n\n");
	END

	START_YLW
#if NORM
	printf("   NORMAL DATA TEST!\n");
#else
	printf("   RANDOM DATA TEST!\n");
#endif
	END

	START_GRN
#if LOOP
	printf("   TEST WAS RUNNED %d TIMES!\n", test_cnt);
#elif SEL
	printf("   TEST WAS RUNNED SAME TIMES!\n");
#else
	printf("   TEST WAS RUNNED ONCE!\n");
#endif
	END

#if REVERSE
	START_BLU
	printf("   TEST WAS RUNNED REVERSED ORDER!\n");
	END
#endif

	for(int i = 0; i < CASE; i++)
	{
		printf("   data size : %s, latency : %ld\n", size_str[i], latency[i]);
	}

	START_RED
	printf("\n______________________TEST END_______________________\n\n\n");
	END
}


int main(void)
{
	int test_cnt;

	char * device_buf;
	char * host_buf;

	srand(time(NULL));

	host_buf = (char *)calloc(HALF * 2, sizeof(char));
	hipHostAlloc((void**)&device_buf, HALF * sizeof(char), hipHostMallocDefault);
	hipMemset(device_buf, 0 ,HALF * sizeof(char)); 

	printf("Enter the test_cnt\n");
	scanf("%d", &test_cnt);

#if LOOP
	int loop_cnt;

	printf("Enter the loop_cnt\n");
	scanf("%d", &loop_cnt);
	loop(device_buf, host_buf, test_cnt, loop_cnt);
#elif SEL
	same_cnt_loop(device_buf, host_buf, test_cnt);
#else
	once(device_buf, host_buf, test_cnt);
#endif
	print_result(test_cnt);

	hipFree(device_buf);

	fclose(data);

	return 0;
}

	.text
	.file	"memcpy_test.hip"
	.globl	_Z14monotonic_timev             # -- Begin function _Z14monotonic_timev
	.type	_Z14monotonic_timev,@function
_Z14monotonic_timev:                    # @_Z14monotonic_timev
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	subq	$16, %rsp
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -16
	movq	%rsp, %rbx
	movl	$1, %edi
	movq	%rbx, %rsi
	callq	clock_gettime
	imull	$1000000000, (%rbx), %eax       # imm = 0x3B9ACA00
	addl	8(%rbx), %eax
	addq	$16, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end0:
	.size	_Z14monotonic_timev, .Lfunc_end0-_Z14monotonic_timev
	.cfi_endproc
                                        # -- End function
	.globl	_Z9call_dataPci                 # -- Begin function _Z9call_dataPci
	.type	_Z9call_dataPci,@function
_Z9call_dataPci:                        # @_Z9call_dataPci
	.cfi_startproc
# %bb.0:
	pushq	%r14
	.cfi_def_cfa_offset 16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	pushq	%rax
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -24
	.cfi_offset %r14, -16
	movl	%esi, %ebx
	movq	%rdi, %r14
	movq	data(%rip), %rdi
	xorl	%esi, %esi
	xorl	%edx, %edx
	callq	fseek
	movq	data(%rip), %rdx
	movq	%r14, %rdi
	movl	%ebx, %esi
	addq	$8, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	jmp	fgets                           # TAILCALL
.Lfunc_end1:
	.size	_Z9call_dataPci, .Lfunc_end1-_Z9call_dataPci
	.cfi_endproc
                                        # -- End function
	.globl	_Z4oncePcS_i                    # -- Begin function _Z4oncePcS_i
	.type	_Z4oncePcS_i,@function
_Z4oncePcS_i:                           # @_Z4oncePcS_i
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	subq	$24, %rsp
	.cfi_def_cfa_offset 80
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movl	%edx, 4(%rsp)                   # 4-byte Spill
	testl	%edx, %edx
	jle	.LBB2_5
# %bb.1:                                # %.lr.ph
	movq	%rsi, %r14
	movq	%rdi, %r15
	xorl	%r13d, %r13d
	leaq	8(%rsp), %r12
.LBB2_2:                                # =>This Loop Header: Depth=1
                                        #     Child Loop BB2_3 Depth 2
	movq	data(%rip), %rdi
	xorl	%esi, %esi
	xorl	%edx, %edx
	callq	fseek
	movq	data(%rip), %rdx
	movq	%r14, %rdi
	movl	$4194304, %esi                  # imm = 0x400000
	callq	fgets
	movq	$-68, %rbp
.LBB2_3:                                #   Parent Loop BB2_2 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	callq	rand
	movl	%eax, %ebx
	andl	$2097151, %ebx                  # imm = 0x1FFFFF
	cmpl	$2097152, %eax                  # imm = 0x200000
	cmovll	%eax, %ebx
	movl	$1, %edi
	movq	%r12, %rsi
	callq	clock_gettime
	imull	$1000000000, 8(%rsp), %eax      # imm = 0x3B9ACA00
	addl	16(%rsp), %eax
	movl	%eax, start+68(%rbp)
	movslq	%ebx, %rsi
	addq	%r14, %rsi
	movslq	size+68(%rbp), %rdx
	movq	%r15, %rdi
	movl	$1, %ecx
	callq	hipMemcpy
	movl	$1, %edi
	movq	%r12, %rsi
	callq	clock_gettime
	imull	$1000000000, 8(%rsp), %eax      # imm = 0x3B9ACA00
	addl	16(%rsp), %eax
	movl	%eax, end+68(%rbp)
	movslq	start+68(%rbp), %rcx
	cltq
	subq	%rcx, %rax
	addq	%rax, latency+136(%rbp,%rbp)
	movslq	size+68(%rbp), %rdx
	movq	%r15, %rdi
	xorl	%esi, %esi
	callq	hipMemset
	addq	$4, %rbp
	jne	.LBB2_3
# %bb.4:                                #   in Loop: Header=BB2_2 Depth=1
	incl	%r13d
	cmpl	4(%rsp), %r13d                  # 4-byte Folded Reload
	jne	.LBB2_2
.LBB2_5:                                # %.preheader
	movslq	4(%rsp), %rcx                   # 4-byte Folded Reload
	movq	$-136, %rsi
.LBB2_6:                                # =>This Inner Loop Header: Depth=1
	movq	latency+136(%rsi), %rax
	xorl	%edx, %edx
	divq	%rcx
	movq	%rax, latency+136(%rsi)
	addq	$8, %rsi
	jne	.LBB2_6
# %bb.7:
	addq	$24, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end2:
	.size	_Z4oncePcS_i, .Lfunc_end2-_Z4oncePcS_i
	.cfi_endproc
                                        # -- End function
	.globl	_Z4loopPcS_ii                   # -- Begin function _Z4loopPcS_ii
	.type	_Z4loopPcS_ii,@function
_Z4loopPcS_ii:                          # @_Z4loopPcS_ii
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	subq	$56, %rsp
	.cfi_def_cfa_offset 112
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movq	%rsi, 16(%rsp)                  # 8-byte Spill
	movl	%edx, 12(%rsp)                  # 4-byte Spill
	testl	%edx, %edx
	jle	.LBB3_8
# %bb.1:                                # %.lr.ph35
	movl	%ecx, %ebx
	movq	%rdi, %r15
	movslq	%ecx, %rax
	movq	%rax, 48(%rsp)                  # 8-byte Spill
	xorl	%eax, %eax
	leaq	24(%rsp), %r12
.LBB3_2:                                # =>This Loop Header: Depth=1
                                        #     Child Loop BB3_3 Depth 2
                                        #       Child Loop BB3_5 Depth 3
	movq	%rax, 40(%rsp)                  # 8-byte Spill
	movq	data(%rip), %rdi
	xorl	%esi, %esi
	xorl	%edx, %edx
	callq	fseek
	movq	data(%rip), %rdx
	movq	16(%rsp), %rdi                  # 8-byte Reload
	movl	$4194304, %esi                  # imm = 0x400000
	callq	fgets
	xorl	%r14d, %r14d
.LBB3_3:                                #   Parent Loop BB3_2 Depth=1
                                        # =>  This Loop Header: Depth=2
                                        #       Child Loop BB3_5 Depth 3
	callq	rand
	movl	%eax, %r13d
	movl	$1, %edi
	movq	%r12, %rsi
	callq	clock_gettime
	imull	$1000000000, 24(%rsp), %eax     # imm = 0x3B9ACA00
	addl	32(%rsp), %eax
	movl	%eax, start(,%r14,4)
	testl	%ebx, %ebx
	jle	.LBB3_6
# %bb.4:                                # %.lr.ph
                                        #   in Loop: Header=BB3_3 Depth=2
	movl	%r13d, %eax
	andl	$2097151, %eax                  # imm = 0x1FFFFF
	cmpl	$2097152, %r13d                 # imm = 0x200000
	cmovll	%r13d, %eax
	movslq	%eax, %r13
	addq	16(%rsp), %r13                  # 8-byte Folded Reload
	movl	%ebx, %ebp
.LBB3_5:                                #   Parent Loop BB3_2 Depth=1
                                        #     Parent Loop BB3_3 Depth=2
                                        # =>    This Inner Loop Header: Depth=3
	movslq	size(,%r14,4), %rdx
	movq	%r15, %rdi
	movq	%r13, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	decl	%ebp
	jne	.LBB3_5
.LBB3_6:                                # %._crit_edge
                                        #   in Loop: Header=BB3_3 Depth=2
	movl	$1, %edi
	movq	%r12, %rsi
	callq	clock_gettime
	imulq	$1000000000, 24(%rsp), %rax     # imm = 0x3B9ACA00
	addq	32(%rsp), %rax
	movl	%eax, end(,%r14,4)
	movslq	start(,%r14,4), %rcx
	cltq
	subq	%rcx, %rax
	xorl	%edx, %edx
	divq	48(%rsp)                        # 8-byte Folded Reload
	addq	%rax, latency(,%r14,8)
	incq	%r14
	cmpq	$17, %r14
	jne	.LBB3_3
# %bb.7:                                #   in Loop: Header=BB3_2 Depth=1
	movq	40(%rsp), %rax                  # 8-byte Reload
	incl	%eax
	cmpl	12(%rsp), %eax                  # 4-byte Folded Reload
	jne	.LBB3_2
.LBB3_8:                                # %.preheader
	movslq	12(%rsp), %rcx                  # 4-byte Folded Reload
	movq	$-136, %rsi
.LBB3_9:                                # =>This Inner Loop Header: Depth=1
	movq	latency+136(%rsi), %rax
	xorl	%edx, %edx
	divq	%rcx
	movq	%rax, latency+136(%rsi)
	addq	$8, %rsi
	jne	.LBB3_9
# %bb.10:
	addq	$56, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end3:
	.size	_Z4loopPcS_ii, .Lfunc_end3-_Z4loopPcS_ii
	.cfi_endproc
                                        # -- End function
	.globl	_Z13same_cnt_loopPcS_i          # -- Begin function _Z13same_cnt_loopPcS_i
	.type	_Z13same_cnt_loopPcS_i,@function
_Z13same_cnt_loopPcS_i:                 # @_Z13same_cnt_loopPcS_i
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	subq	$40, %rsp
	.cfi_def_cfa_offset 96
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movq	%rsi, 8(%rsp)                   # 8-byte Spill
	movl	%edx, 4(%rsp)                   # 4-byte Spill
	testl	%edx, %edx
	jle	.LBB4_8
# %bb.1:                                # %.lr.ph35
	movq	%rdi, %r15
	movl	size+64(%rip), %ebp
	xorl	%eax, %eax
	leaq	16(%rsp), %r12
.LBB4_2:                                # =>This Loop Header: Depth=1
                                        #     Child Loop BB4_3 Depth 2
                                        #       Child Loop BB4_5 Depth 3
	movq	%rax, 32(%rsp)                  # 8-byte Spill
	movq	data(%rip), %rdi
	xorl	%esi, %esi
	xorl	%edx, %edx
	callq	fseek
	movq	data(%rip), %rdx
	movq	8(%rsp), %rdi                   # 8-byte Reload
	movl	$4194304, %esi                  # imm = 0x400000
	callq	fgets
	xorl	%r14d, %r14d
.LBB4_3:                                #   Parent Loop BB4_2 Depth=1
                                        # =>  This Loop Header: Depth=2
                                        #       Child Loop BB4_5 Depth 3
	movl	%ebp, %eax
	cltd
	idivl	size(,%r14,4)
	movl	%eax, %r13d
	callq	rand
	movl	%eax, %ebx
	movl	$1, %edi
	movq	%r12, %rsi
	callq	clock_gettime
	imull	$1000000000, 16(%rsp), %eax     # imm = 0x3B9ACA00
	addl	24(%rsp), %eax
	movl	%eax, start(,%r14,4)
	testl	%r13d, %r13d
	jle	.LBB4_6
# %bb.4:                                # %.lr.ph
                                        #   in Loop: Header=BB4_3 Depth=2
	movl	%ebx, %eax
	andl	$2097151, %eax                  # imm = 0x1FFFFF
	cmpl	$2097152, %ebx                  # imm = 0x200000
	cmovll	%ebx, %eax
	movslq	%eax, %rbx
	addq	8(%rsp), %rbx                   # 8-byte Folded Reload
.LBB4_5:                                #   Parent Loop BB4_2 Depth=1
                                        #     Parent Loop BB4_3 Depth=2
                                        # =>    This Inner Loop Header: Depth=3
	movslq	size(,%r14,4), %rdx
	movq	%r15, %rdi
	movq	%rbx, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	decl	%r13d
	jne	.LBB4_5
.LBB4_6:                                # %._crit_edge
                                        #   in Loop: Header=BB4_3 Depth=2
	movl	$1, %edi
	movq	%r12, %rsi
	callq	clock_gettime
	imull	$1000000000, 16(%rsp), %eax     # imm = 0x3B9ACA00
	addl	24(%rsp), %eax
	movslq	start(,%r14,4), %rcx
	movslq	%eax, %rdx
	subq	%rcx, %rdx
	addq	%rdx, latency(,%r14,8)
	movl	%eax, end(,%r14,4)
	incq	%r14
	cmpq	$17, %r14
	jne	.LBB4_3
# %bb.7:                                #   in Loop: Header=BB4_2 Depth=1
	movq	32(%rsp), %rax                  # 8-byte Reload
	incl	%eax
	cmpl	4(%rsp), %eax                   # 4-byte Folded Reload
	jne	.LBB4_2
.LBB4_8:                                # %.preheader
	movslq	4(%rsp), %rcx                   # 4-byte Folded Reload
	movq	$-136, %rsi
.LBB4_9:                                # =>This Inner Loop Header: Depth=1
	movq	latency+136(%rsi), %rax
	xorl	%edx, %edx
	divq	%rcx
	movq	%rax, latency+136(%rsi)
	addq	$8, %rsi
	jne	.LBB4_9
# %bb.10:
	addq	$40, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end4:
	.size	_Z13same_cnt_loopPcS_i, .Lfunc_end4-_Z13same_cnt_loopPcS_i
	.cfi_endproc
                                        # -- End function
	.globl	_Z12print_resulti               # -- Begin function _Z12print_resulti
	.type	_Z12print_resulti,@function
_Z12print_resulti:                      # @_Z12print_resulti
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset %rbx, -16
	movl	$.L.str.19, %edi
	xorl	%eax, %eax
	callq	printf
	movl	$.Lstr, %edi
	callq	puts@PLT
	movl	$.L.str.21, %edi
	xorl	%eax, %eax
	callq	printf
	movl	$.L.str.22, %edi
	xorl	%eax, %eax
	callq	printf
	movl	$.Lstr.1, %edi
	callq	puts@PLT
	movl	$.L.str.21, %edi
	xorl	%eax, %eax
	callq	printf
	movl	$.L.str.24, %edi
	xorl	%eax, %eax
	callq	printf
	movl	$.Lstr.2, %edi
	callq	puts@PLT
	movl	$.L.str.21, %edi
	xorl	%eax, %eax
	callq	printf
	movq	$-136, %rbx
.LBB5_1:                                # =>This Inner Loop Header: Depth=1
	movq	size_str+136(%rbx), %rsi
	movq	latency+136(%rbx), %rdx
	movl	$.L.str.26, %edi
	xorl	%eax, %eax
	callq	printf
	addq	$8, %rbx
	jne	.LBB5_1
# %bb.2:
	movl	$.L.str.19, %edi
	xorl	%eax, %eax
	callq	printf
	movl	$.Lstr.3, %edi
	callq	puts@PLT
	movl	$.L.str.21, %edi
	xorl	%eax, %eax
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	printf                          # TAILCALL
.Lfunc_end5:
	.size	_Z12print_resulti, .Lfunc_end5-_Z12print_resulti
	.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	%rbx
	.cfi_def_cfa_offset 32
	subq	$16, %rsp
	.cfi_def_cfa_offset 48
	.cfi_offset %rbx, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	xorl	%edi, %edi
	callq	time
	movl	%eax, %edi
	callq	srand
	movl	$4194304, %edi                  # imm = 0x400000
	movl	$1, %esi
	callq	calloc
	movq	%rax, %r14
	leaq	8(%rsp), %rbx
	movl	$2097152, %esi                  # imm = 0x200000
	movq	%rbx, %rdi
	xorl	%edx, %edx
	callq	hipHostAlloc
	movq	(%rbx), %rdi
	movl	$2097152, %edx                  # imm = 0x200000
	xorl	%esi, %esi
	callq	hipMemset
	movl	$.Lstr.4, %edi
	callq	puts@PLT
	leaq	4(%rsp), %r15
	movl	$.L.str.29, %edi
	movq	%r15, %rsi
	xorl	%eax, %eax
	callq	__isoc23_scanf
	movq	(%rbx), %rdi
	movl	(%r15), %edx
	movq	%r14, %rsi
	callq	_Z13same_cnt_loopPcS_i
	callq	_Z12print_resulti
	movq	(%rbx), %rdi
	callq	hipFree
	movq	data(%rip), %rdi
	callq	fclose
	xorl	%eax, %eax
	addq	$16, %rsp
	.cfi_def_cfa_offset 32
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end6:
	.size	main, .Lfunc_end6-main
	.cfi_endproc
                                        # -- End function
	.section	.text.startup,"ax",@progbits
	.type	_GLOBAL__sub_I_memcpy_test.hip,@function # -- Begin function _GLOBAL__sub_I_memcpy_test.hip
_GLOBAL__sub_I_memcpy_test.hip:         # @_GLOBAL__sub_I_memcpy_test.hip
	.cfi_startproc
# %bb.0:
	pushq	%rax
	.cfi_def_cfa_offset 16
	movl	$.L.str, %edi
	movl	$.L.str.1, %esi
	callq	fopen
	movq	%rax, data(%rip)
	popq	%rax
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end7:
	.size	_GLOBAL__sub_I_memcpy_test.hip, .Lfunc_end7-_GLOBAL__sub_I_memcpy_test.hip
	.cfi_endproc
                                        # -- End function
	.type	data,@object                    # @data
	.bss
	.globl	data
	.p2align	3, 0x0
data:
	.quad	0
	.size	data, 8

	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"data.txt"
	.size	.L.str, 9

	.type	.L.str.1,@object                # @.str.1
.L.str.1:
	.asciz	"r"
	.size	.L.str.1, 2

	.type	latency,@object                 # @latency
	.bss
	.globl	latency
	.p2align	4, 0x0
latency:
	.zero	136
	.size	latency, 136

	.type	.L.str.2,@object                # @.str.2
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str.2:
	.asciz	"64"
	.size	.L.str.2, 3

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"128"
	.size	.L.str.3, 4

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

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

	.type	.L.str.6,@object                # @.str.6
.L.str.6:
	.asciz	"1024"
	.size	.L.str.6, 5

	.type	.L.str.7,@object                # @.str.7
.L.str.7:
	.asciz	"1514"
	.size	.L.str.7, 5

	.type	.L.str.8,@object                # @.str.8
.L.str.8:
	.asciz	"64 * 32"
	.size	.L.str.8, 8

	.type	.L.str.9,@object                # @.str.9
.L.str.9:
	.asciz	"64 * 64"
	.size	.L.str.9, 8

	.type	.L.str.10,@object               # @.str.10
.L.str.10:
	.asciz	"64 * 128"
	.size	.L.str.10, 9

	.type	.L.str.11,@object               # @.str.11
.L.str.11:
	.asciz	"64 * 256"
	.size	.L.str.11, 9

	.type	.L.str.12,@object               # @.str.12
.L.str.12:
	.asciz	"64 * 512"
	.size	.L.str.12, 9

	.type	.L.str.13,@object               # @.str.13
.L.str.13:
	.asciz	"64 * 1024"
	.size	.L.str.13, 10

	.type	.L.str.14,@object               # @.str.14
.L.str.14:
	.asciz	"64 * 1024 * 2"
	.size	.L.str.14, 14

	.type	.L.str.15,@object               # @.str.15
.L.str.15:
	.asciz	"64 * 1024 * 4"
	.size	.L.str.15, 14

	.type	.L.str.16,@object               # @.str.16
.L.str.16:
	.asciz	"64 * 1024 * 8"
	.size	.L.str.16, 14

	.type	.L.str.17,@object               # @.str.17
.L.str.17:
	.asciz	"64 * 1024 * 16"
	.size	.L.str.17, 15

	.type	.L.str.18,@object               # @.str.18
.L.str.18:
	.asciz	"64 * 1024 * 32"
	.size	.L.str.18, 15

	.type	size_str,@object                # @size_str
	.data
	.globl	size_str
	.p2align	4, 0x0
size_str:
	.quad	.L.str.2
	.quad	.L.str.3
	.quad	.L.str.4
	.quad	.L.str.5
	.quad	.L.str.6
	.quad	.L.str.7
	.quad	.L.str.8
	.quad	.L.str.9
	.quad	.L.str.10
	.quad	.L.str.11
	.quad	.L.str.12
	.quad	.L.str.13
	.quad	.L.str.14
	.quad	.L.str.15
	.quad	.L.str.16
	.quad	.L.str.17
	.quad	.L.str.18
	.size	size_str, 136

	.type	size,@object                    # @size
	.globl	size
	.p2align	4, 0x0
size:
	.long	64                              # 0x40
	.long	128                             # 0x80
	.long	256                             # 0x100
	.long	512                             # 0x200
	.long	1024                            # 0x400
	.long	1514                            # 0x5ea
	.long	2048                            # 0x800
	.long	4096                            # 0x1000
	.long	8192                            # 0x2000
	.long	16384                           # 0x4000
	.long	32768                           # 0x8000
	.long	65536                           # 0x10000
	.long	131072                          # 0x20000
	.long	262144                          # 0x40000
	.long	524288                          # 0x80000
	.long	1048576                         # 0x100000
	.long	2097152                         # 0x200000
	.size	size, 68

	.type	start,@object                   # @start
	.bss
	.globl	start
	.p2align	4, 0x0
start:
	.zero	68
	.size	start, 68

	.type	end,@object                     # @end
	.globl	end
	.p2align	4, 0x0
end:
	.zero	68
	.size	end, 68

	.type	.L.str.19,@object               # @.str.19
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str.19:
	.asciz	"\033[1;31m"
	.size	.L.str.19, 8

	.type	.L.str.21,@object               # @.str.21
.L.str.21:
	.asciz	"\033[0m"
	.size	.L.str.21, 5

	.type	.L.str.22,@object               # @.str.22
.L.str.22:
	.asciz	"\033[1;33m"
	.size	.L.str.22, 8

	.type	.L.str.24,@object               # @.str.24
.L.str.24:
	.asciz	"\033[1;32m"
	.size	.L.str.24, 8

	.type	.L.str.26,@object               # @.str.26
.L.str.26:
	.asciz	"   data size : %s, latency : %ld\n"
	.size	.L.str.26, 34

	.type	.L.str.29,@object               # @.str.29
.L.str.29:
	.asciz	"%d"
	.size	.L.str.29, 3

	.section	.init_array,"aw",@init_array
	.p2align	3, 0x0
	.quad	_GLOBAL__sub_I_memcpy_test.hip
	.type	__hip_cuid_,@object             # @__hip_cuid_
	.bss
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               # 0x0
	.size	__hip_cuid_, 1

	.type	.Lstr,@object                   # @str
	.section	.rodata.str1.1,"aMS",@progbits,1
.Lstr:
	.asciz	"\n\n______________________TEST START_______________________\n"
	.size	.Lstr, 59

	.type	.Lstr.1,@object                 # @str.1
.Lstr.1:
	.asciz	"   RANDOM DATA TEST!"
	.size	.Lstr.1, 21

	.type	.Lstr.2,@object                 # @str.2
.Lstr.2:
	.asciz	"   TEST WAS RUNNED SAME TIMES!"
	.size	.Lstr.2, 31

	.type	.Lstr.3,@object                 # @str.3
.Lstr.3:
	.asciz	"\n______________________TEST END_______________________\n\n"
	.size	.Lstr.3, 57

	.type	.Lstr.4,@object                 # @str.4
.Lstr.4:
	.asciz	"Enter the test_cnt"
	.size	.Lstr.4, 19

	.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 _GLOBAL__sub_I_memcpy_test.hip
	.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
