#include <iostream>
#include <chrono>
#include <cuda_runtime.h>
#include <string>
#include <iomanip>

using namespace std;

using ST = unsigned long long;

constexpr ST TOTAL_SIZE = 1 << 30; // 1 GB
constexpr ST TOTAL_SIZE_IN_BYTES = TOTAL_SIZE * sizeof(char);
constexpr ST CNT = 19;

const string grand_name[CNT] = {
	"4KB", "8KB", "16KB", "32KB", "64KB", "128KB", "256KB", "512KB", "1MB", "2MB",
	"4MB", "8MB", "16MB", "32MB", "64MB", "128MB", "256MB", "512MB", "1GB"
};

const ST grand_size[CNT] = {
	1 << 12, 1 << 13, 1 << 14, 1 << 15,
	1 << 16, 1 << 17, 1 << 18, 1 << 19,
	1 << 20, 1 << 21, 1 << 22, 1 << 23,
	1 << 24, 1 << 25, 1 << 26, 1 << 27,
	1 << 28, 1 << 29, 1 << 30
};



// granularity

double copyTest(char *dst, char *src, ST total, ST granularity, cudaMemcpyKind kind) {
	ST idx = 0;
	auto t1 = chrono::system_clock::now();
	for (; idx < total; idx += granularity) {
		cudaMemcpy(dst + idx, src + idx, granularity * sizeof(char), kind);
	}
	auto t2 = chrono::system_clock::now();
	chrono::duration<double> diff = t2 - t1;
	return diff.count();
}

void rangeTest(char *dst, char *src, ST total, cudaMemcpyKind kind) {
	for (int i = 0; i < CNT; i++) {
		double tm = copyTest(dst, src, total, grand_size[i], kind);
		double tp = 1024 / tm; // MB/s
		cout << fixed << tp << endl;
	}
}


int main() {
	char *buf = new char[TOTAL_SIZE];
	char *buf2 = new char[TOTAL_SIZE];
	char *gpu_buf, *gpu_buf2;
	cout << "granularity: " << endl;
	for (const auto &s : grand_name)
		cout << s << endl;
	cout << endl;

	cudaMalloc((void **)&gpu_buf, TOTAL_SIZE_IN_BYTES);
	cudaMalloc((void **)&gpu_buf2, TOTAL_SIZE_IN_BYTES);
	memset(buf, 'a', TOTAL_SIZE);
	cout << "host to device: " << endl;
	rangeTest(gpu_buf, buf, TOTAL_SIZE, cudaMemcpyHostToDevice);
	cout << endl;
	cout << "(MB/s)" << endl;

	cout << "device to device: " << endl;
	rangeTest(gpu_buf, gpu_buf2, TOTAL_SIZE, cudaMemcpyDeviceToDevice);
	cout << endl;
	cout << "(MB/s)" << endl;

	cout << "device to host: " << endl;
	rangeTest(buf2, gpu_buf2, TOTAL_SIZE, cudaMemcpyDeviceToHost);
	cout << endl;
	cout << buf2[32];
	cout << "(MB/s)" << endl;

	cudaFree(gpu_buf);
	cudaFree(gpu_buf2);
	delete[] buf;
	delete[] buf2;


//#ifdef WIN32
//	system("pause");
//#endif

	return 0;
}
	.file	"tmpxft_00307d6c_00000000-6_BandwidthTrail.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

#NO_APP
	.section	.text._ZSt5fixedRSt8ios_base,"axG",@progbits,_ZSt5fixedRSt8ios_base,comdat
	.weak	_ZSt5fixedRSt8ios_base
	.type	_ZSt5fixedRSt8ios_base, @function
_ZSt5fixedRSt8ios_base:
.LFB3285:
	.cfi_startproc
	endbr64
	movl	24(%rdi), %edx
	movq	%rdi, %rax
	andl	$-261, %edx
	orl	$4, %edx
	movl	%edx, 24(%rdi)
	ret
	.cfi_endproc
.LFE3285:
	.size	_ZSt5fixedRSt8ios_base, .-_ZSt5fixedRSt8ios_base
	.text
	.type	__tcf_0, @function
__tcf_0:
.LFB4748:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	leaq	608+_ZL10grand_name(%rip), %rbx
	pushq	%rdx
	.cfi_def_cfa_offset 32
	leaq	-608(%rbx), %rbp
.L3:
	subq	$32, %rbx
	movq	%rbx, %rdi
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE10_M_disposeEv@PLT
	cmpq	%rbp, %rbx
	jne	.L3
	popq	%rax
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4748:
	.size	__tcf_0, .-__tcf_0
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB4023:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE4023:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"basic_string: construction from null is not valid"
	.section	.text.startup,"ax",@progbits
	.align 2
	.type	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_.constprop.0, @function
_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_.constprop.0:
.LFB4753:
	.cfi_startproc
	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
	subq	$16, %rsp
	.cfi_def_cfa_offset 48
	movq	%fs:40, %rax
	movq	%rax, 8(%rsp)
	xorl	%eax, %eax
	leaq	16(%rdi), %rax
	movq	%rax, (%rdi)
	testq	%rsi, %rsi
	jne	.L9
	movq	8(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L14
	leaq	.LC0(%rip), %rdi
	call	_ZSt19__throw_logic_errorPKc@PLT
.L9:
	movq	%rdi, %rbx
	movq	%rsi, %rdi
	movq	%rsi, %rbp
	call	strlen@PLT
	movq	%rax, (%rsp)
	leaq	0(%rbp,%rax), %r12
	cmpq	$15, %rax
	jbe	.L11
	movq	%rsp, %rsi
	xorl	%edx, %edx
	movq	%rbx, %rdi
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE9_M_createERmm@PLT
	movq	%rax, (%rbx)
	movq	(%rsp), %rax
	movq	%rax, 16(%rbx)
.L11:
	movq	(%rbx), %rdi
	movq	%r12, %rdx
	movq	%rbp, %rsi
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE13_S_copy_charsEPcPKcS7_@PLT
	movq	(%rsp), %rax
	movq	(%rbx), %rdx
	movq	%rax, 8(%rbx)
	movb	$0, (%rdx,%rax)
	movq	8(%rsp), %rax
	subq	%fs:40, %rax
	je	.L12
.L14:
	call	__stack_chk_fail@PLT
.L12:
	addq	$16, %rsp
	.cfi_def_cfa_offset 32
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%rbp
	.cfi_def_cfa_offset 16
	popq	%r12
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4753:
	.size	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_.constprop.0, .-_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_.constprop.0
	.text
	.align 2
	.type	_ZNSolsEPFRSoS_E.isra.0, @function
_ZNSolsEPFRSoS_E.isra.0:
.LFB4755:
	.cfi_startproc
	jmp	*%rsi
	.cfi_endproc
.LFE4755:
	.size	_ZNSolsEPFRSoS_E.isra.0, .-_ZNSolsEPFRSoS_E.isra.0
	.globl	_Z8copyTestPcS_yy14cudaMemcpyKind
	.type	_Z8copyTestPcS_yy14cudaMemcpyKind, @function
_Z8copyTestPcS_yy14cudaMemcpyKind:
.LFB4014:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	xorl	%r15d, %r15d
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	movq	%rsi, %r14
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	movq	%rdi, %r13
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movq	%rdx, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movq	%rcx, %rbx
	subq	$24, %rsp
	.cfi_def_cfa_offset 80
	movl	%r8d, 12(%rsp)
	call	_ZNSt6chrono3_V212system_clock3nowEv@PLT
	movq	%rax, %rbp
.L17:
	cmpq	%r12, %r15
	jnb	.L20
	movl	12(%rsp), %ecx
	leaq	(%r14,%r15), %rsi
	leaq	0(%r13,%r15), %rdi
	movq	%rbx, %rdx
	addq	%rbx, %r15
	call	cudaMemcpy@PLT
	jmp	.L17
.L20:
	call	_ZNSt6chrono3_V212system_clock3nowEv@PLT
	subq	%rbp, %rax
	cvtsi2sdq	%rax, %xmm0
	divsd	.LC1(%rip), %xmm0
	addq	$24, %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
.LFE4014:
	.size	_Z8copyTestPcS_yy14cudaMemcpyKind, .-_Z8copyTestPcS_yy14cudaMemcpyKind
	.globl	_Z9rangeTestPcS_y14cudaMemcpyKind
	.type	_Z9rangeTestPcS_y14cudaMemcpyKind, @function
_Z9rangeTestPcS_y14cudaMemcpyKind:
.LFB4019:
	.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
	movl	%ecx, %r14d
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	movq	%rdx, %r13
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movq	%rsi, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movq	%rdi, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	leaq	_ZL10grand_size(%rip), %rbx
	leaq	152(%rbx), %r15
	subq	$24, %rsp
	.cfi_def_cfa_offset 80
.L22:
	movq	(%rbx), %rcx
	movl	%r14d, %r8d
	movq	%r13, %rdx
	movq	%r12, %rsi
	movq	%rbp, %rdi
	addq	$8, %rbx
	call	_Z8copyTestPcS_yy14cudaMemcpyKind
	movsd	.LC2(%rip), %xmm1
	leaq	_ZSt5fixedRSt8ios_base(%rip), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	divsd	%xmm0, %xmm1
	movsd	%xmm1, 8(%rsp)
	call	_ZNSolsEPFRSt8ios_baseS0_E@PLT
	movsd	8(%rsp), %xmm0
	movq	%rax, %rdi
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	cmpq	%r15, %rbx
	jne	.L22
	addq	$24, %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
.LFE4019:
	.size	_Z9rangeTestPcS_y14cudaMemcpyKind, .-_Z9rangeTestPcS_y14cudaMemcpyKind
	.section	.rodata.str1.1
.LC3:
	.string	"granularity: "
.LC4:
	.string	"host to device: "
.LC5:
	.string	"(MB/s)"
.LC6:
	.string	"device to device: "
.LC7:
	.string	"device to host: "
	.section	.text.startup
	.globl	main
	.type	main, @function
main:
.LFB4020:
	.cfi_startproc
	endbr64
	pushq	%r14
	.cfi_def_cfa_offset 16
	.cfi_offset 14, -16
	movl	$1073741824, %edi
	pushq	%r13
	.cfi_def_cfa_offset 24
	.cfi_offset 13, -24
	leaq	_ZL10grand_name(%rip), %r13
	pushq	%r12
	.cfi_def_cfa_offset 32
	.cfi_offset 12, -32
	leaq	608(%r13), %r14
	pushq	%rbp
	.cfi_def_cfa_offset 40
	.cfi_offset 6, -40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset 3, -48
	leaq	_ZSt4cout(%rip), %rbx
	subq	$32, %rsp
	.cfi_def_cfa_offset 80
	movq	%fs:40, %rax
	movq	%rax, 24(%rsp)
	xorl	%eax, %eax
	call	_Znam@PLT
	movl	$1073741824, %edi
	movq	%rax, %rbp
	call	_Znam@PLT
	leaq	.LC3(%rip), %rsi
	movq	%rbx, %rdi
	movq	%rax, %r12
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
.L26:
	movq	%r13, %rsi
	movq	%rbx, %rdi
	addq	$32, %r13
	call	_ZStlsIcSt11char_traitsIcESaIcEERSt13basic_ostreamIT_T0_ES7_RKNSt7__cxx1112basic_stringIS4_S5_T1_EE@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	cmpq	%r14, %r13
	jne	.L26
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rbx, %rdi
	leaq	.LC5(%rip), %r13
	call	_ZNSolsEPFRSoS_E.isra.0
	leaq	8(%rsp), %rdi
	movl	$1073741824, %esi
	call	cudaMalloc@PLT
	leaq	16(%rsp), %rdi
	movl	$1073741824, %esi
	call	cudaMalloc@PLT
	movb	$97, %al
	movl	$1073741824, %ecx
	movq	%rbp, %rdi
	rep stosb
	leaq	.LC4(%rip), %rsi
	movq	%rbx, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	movq	8(%rsp), %rdi
	movl	$1, %ecx
	movq	%rbp, %rsi
	movl	$1073741824, %edx
	call	_Z9rangeTestPcS_y14cudaMemcpyKind
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rbx, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	movq	%r13, %rsi
	movq	%rbx, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	leaq	.LC6(%rip), %rsi
	movq	%rbx, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	movq	16(%rsp), %rsi
	movq	8(%rsp), %rdi
	movl	$3, %ecx
	movl	$1073741824, %edx
	call	_Z9rangeTestPcS_y14cudaMemcpyKind
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rbx, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	movq	%r13, %rsi
	movq	%rbx, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	leaq	.LC7(%rip), %rsi
	movq	%rbx, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	movq	16(%rsp), %rsi
	movl	$2, %ecx
	movq	%r12, %rdi
	movl	$1073741824, %edx
	call	_Z9rangeTestPcS_y14cudaMemcpyKind
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rbx, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	movq	_ZSt4cout(%rip), %rax
	movsbl	32(%r12), %esi
	movq	-24(%rax), %rax
	movb	%sil, 7(%rsp)
	cmpq	$0, 16(%rbx,%rax)
	je	.L27
	leaq	7(%rsp), %rsi
	movl	$1, %edx
	movq	%rbx, %rdi
	call	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l@PLT
	jmp	.L28
.L27:
	movq	%rbx, %rdi
	call	_ZNSo3putEc@PLT
.L28:
	movq	%r13, %rsi
	movq	%rbx, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	movq	8(%rsp), %rdi
	call	cudaFree@PLT
	movq	16(%rsp), %rdi
	call	cudaFree@PLT
	movq	%rbp, %rdi
	call	_ZdaPv@PLT
	movq	%r12, %rdi
	call	_ZdaPv@PLT
	movq	24(%rsp), %rax
	subq	%fs:40, %rax
	je	.L29
	call	__stack_chk_fail@PLT
.L29:
	addq	$32, %rsp
	.cfi_def_cfa_offset 48
	xorl	%eax, %eax
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%rbp
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r13
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4020:
	.size	main, .-main
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB4046:
	.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
.LFE4046:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.rodata.str1.1
.LC8:
	.string	"4KB"
.LC9:
	.string	"8KB"
.LC10:
	.string	"16KB"
.LC11:
	.string	"32KB"
.LC12:
	.string	"64KB"
.LC13:
	.string	"128KB"
.LC14:
	.string	"256KB"
.LC15:
	.string	"512KB"
.LC16:
	.string	"1MB"
.LC17:
	.string	"2MB"
.LC18:
	.string	"4MB"
.LC19:
	.string	"8MB"
.LC20:
	.string	"16MB"
.LC21:
	.string	"32MB"
.LC22:
	.string	"64MB"
.LC23:
	.string	"128MB"
.LC24:
	.string	"256MB"
.LC25:
	.string	"512MB"
.LC26:
	.string	"1GB"
	.section	.text.startup
	.type	_GLOBAL__sub_I__Z8copyTestPcS_yy14cudaMemcpyKind, @function
_GLOBAL__sub_I__Z8copyTestPcS_yy14cudaMemcpyKind:
.LFB4749:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA4749
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	leaq	.LC8(%rip), %rsi
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	leaq	_ZL10grand_name(%rip), %rbx
	movq	%rbx, %rdi
	subq	$24, %rsp
	.cfi_def_cfa_offset 48
.LEHB0:
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_.constprop.0
.LEHE0:
	leaq	.LC9(%rip), %rsi
	leaq	32(%rbx), %rdi
.LEHB1:
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_.constprop.0
.LEHE1:
	jmp	.L113
.L77:
	endbr64
.L36:
	movl	$17, %edx
	jmp	.L37
.L113:
	leaq	.LC10(%rip), %rsi
	leaq	64+_ZL10grand_name(%rip), %rdi
.LEHB2:
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_.constprop.0
.LEHE2:
	jmp	.L114
.L78:
	endbr64
.L39:
	movl	$16, %edx
	jmp	.L37
.L114:
	leaq	.LC11(%rip), %rsi
	leaq	96+_ZL10grand_name(%rip), %rdi
.LEHB3:
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_.constprop.0
.LEHE3:
	jmp	.L115
.L79:
	endbr64
.L41:
	movl	$15, %edx
	jmp	.L37
.L115:
	leaq	.LC12(%rip), %rsi
	leaq	128+_ZL10grand_name(%rip), %rdi
.LEHB4:
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_.constprop.0
.LEHE4:
	jmp	.L116
.L80:
	endbr64
.L43:
	movl	$14, %edx
	jmp	.L37
.L116:
	leaq	.LC13(%rip), %rsi
	leaq	160+_ZL10grand_name(%rip), %rdi
.LEHB5:
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_.constprop.0
.LEHE5:
	jmp	.L117
.L81:
	endbr64
.L45:
	movl	$13, %edx
	jmp	.L37
.L117:
	leaq	.LC14(%rip), %rsi
	leaq	192+_ZL10grand_name(%rip), %rdi
.LEHB6:
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_.constprop.0
.LEHE6:
	jmp	.L118
.L82:
	endbr64
.L47:
	movl	$12, %edx
	jmp	.L37
.L118:
	leaq	.LC15(%rip), %rsi
	leaq	224+_ZL10grand_name(%rip), %rdi
.LEHB7:
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_.constprop.0
.LEHE7:
	jmp	.L119
.L83:
	endbr64
.L49:
	movl	$11, %edx
	jmp	.L37
.L119:
	leaq	.LC16(%rip), %rsi
	leaq	256+_ZL10grand_name(%rip), %rdi
.LEHB8:
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_.constprop.0
.LEHE8:
	jmp	.L120
.L84:
	endbr64
.L51:
	movl	$10, %edx
	jmp	.L37
.L120:
	leaq	.LC17(%rip), %rsi
	leaq	288+_ZL10grand_name(%rip), %rdi
.LEHB9:
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_.constprop.0
.LEHE9:
	jmp	.L121
.L85:
	endbr64
.L53:
	movl	$9, %edx
	jmp	.L37
.L121:
	leaq	.LC18(%rip), %rsi
	leaq	320+_ZL10grand_name(%rip), %rdi
.LEHB10:
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_.constprop.0
.LEHE10:
	jmp	.L122
.L86:
	endbr64
.L55:
	movl	$8, %edx
	jmp	.L37
.L122:
	leaq	.LC19(%rip), %rsi
	leaq	352+_ZL10grand_name(%rip), %rdi
.LEHB11:
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_.constprop.0
.LEHE11:
	jmp	.L123
.L87:
	endbr64
.L57:
	movl	$7, %edx
	jmp	.L37
.L123:
	leaq	.LC20(%rip), %rsi
	leaq	384+_ZL10grand_name(%rip), %rdi
.LEHB12:
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_.constprop.0
.LEHE12:
	jmp	.L124
.L88:
	endbr64
.L59:
	movl	$6, %edx
	jmp	.L37
.L124:
	leaq	.LC21(%rip), %rsi
	leaq	416+_ZL10grand_name(%rip), %rdi
.LEHB13:
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_.constprop.0
.LEHE13:
	jmp	.L125
.L89:
	endbr64
.L61:
	movl	$5, %edx
	jmp	.L37
.L125:
	leaq	.LC22(%rip), %rsi
	leaq	448+_ZL10grand_name(%rip), %rdi
.LEHB14:
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_.constprop.0
.LEHE14:
	jmp	.L126
.L90:
	endbr64
.L63:
	movl	$4, %edx
	jmp	.L37
.L126:
	leaq	.LC23(%rip), %rsi
	leaq	480+_ZL10grand_name(%rip), %rdi
.LEHB15:
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_.constprop.0
.LEHE15:
	jmp	.L127
.L91:
	endbr64
.L65:
	movl	$3, %edx
	jmp	.L37
.L127:
	leaq	.LC24(%rip), %rsi
	leaq	512+_ZL10grand_name(%rip), %rdi
.LEHB16:
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_.constprop.0
.LEHE16:
	jmp	.L128
.L92:
	endbr64
.L67:
	movl	$2, %edx
	jmp	.L37
.L128:
	leaq	.LC25(%rip), %rsi
	leaq	544+_ZL10grand_name(%rip), %rdi
.LEHB17:
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_.constprop.0
.LEHE17:
	jmp	.L129
.L93:
	endbr64
.L69:
	movl	$1, %edx
	jmp	.L37
.L129:
	leaq	.LC26(%rip), %rsi
	leaq	576+_ZL10grand_name(%rip), %rdi
.LEHB18:
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEC2IS3_EEPKcRKS3_.constprop.0
.LEHE18:
	addq	$24, %rsp
	.cfi_remember_state
	.cfi_def_cfa_offset 24
	leaq	__dso_handle(%rip), %rdx
	xorl	%esi, %esi
	popq	%rbx
	.cfi_def_cfa_offset 16
	leaq	__tcf_0(%rip), %rdi
	popq	%rbp
	.cfi_def_cfa_offset 8
	jmp	__cxa_atexit@PLT
.L75:
	.cfi_restore_state
	endbr64
.L71:
	xorl	%edx, %edx
	jmp	.L37
.L76:
	endbr64
.L72:
	movl	$18, %edx
.L37:
	movl	$18, %ebp
	subq	%rdx, %rbp
	salq	$5, %rbp
	addq	%rbx, %rbp
.L74:
	cmpq	%rbx, %rbp
	je	.L73
	subq	$32, %rbp
	movq	%rax, 8(%rsp)
	movq	%rbp, %rdi
	call	_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE10_M_disposeEv@PLT
	movq	8(%rsp), %rax
	jmp	.L74
.L73:
	movq	%rax, %rdi
.LEHB19:
	call	_Unwind_Resume@PLT
.LEHE19:
	.cfi_endproc
.LFE4749:
	.globl	__gxx_personality_v0
	.section	.gcc_except_table,"a",@progbits
.LLSDA4749:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE4749-.LLSDACSB4749
.LLSDACSB4749:
	.uleb128 .LEHB0-.LFB4749
	.uleb128 .LEHE0-.LEHB0
	.uleb128 .L76-.LFB4749
	.uleb128 0
	.uleb128 .LEHB1-.LFB4749
	.uleb128 .LEHE1-.LEHB1
	.uleb128 .L77-.LFB4749
	.uleb128 0
	.uleb128 .LEHB2-.LFB4749
	.uleb128 .LEHE2-.LEHB2
	.uleb128 .L78-.LFB4749
	.uleb128 0
	.uleb128 .LEHB3-.LFB4749
	.uleb128 .LEHE3-.LEHB3
	.uleb128 .L79-.LFB4749
	.uleb128 0
	.uleb128 .LEHB4-.LFB4749
	.uleb128 .LEHE4-.LEHB4
	.uleb128 .L80-.LFB4749
	.uleb128 0
	.uleb128 .LEHB5-.LFB4749
	.uleb128 .LEHE5-.LEHB5
	.uleb128 .L81-.LFB4749
	.uleb128 0
	.uleb128 .LEHB6-.LFB4749
	.uleb128 .LEHE6-.LEHB6
	.uleb128 .L82-.LFB4749
	.uleb128 0
	.uleb128 .LEHB7-.LFB4749
	.uleb128 .LEHE7-.LEHB7
	.uleb128 .L83-.LFB4749
	.uleb128 0
	.uleb128 .LEHB8-.LFB4749
	.uleb128 .LEHE8-.LEHB8
	.uleb128 .L84-.LFB4749
	.uleb128 0
	.uleb128 .LEHB9-.LFB4749
	.uleb128 .LEHE9-.LEHB9
	.uleb128 .L85-.LFB4749
	.uleb128 0
	.uleb128 .LEHB10-.LFB4749
	.uleb128 .LEHE10-.LEHB10
	.uleb128 .L86-.LFB4749
	.uleb128 0
	.uleb128 .LEHB11-.LFB4749
	.uleb128 .LEHE11-.LEHB11
	.uleb128 .L87-.LFB4749
	.uleb128 0
	.uleb128 .LEHB12-.LFB4749
	.uleb128 .LEHE12-.LEHB12
	.uleb128 .L88-.LFB4749
	.uleb128 0
	.uleb128 .LEHB13-.LFB4749
	.uleb128 .LEHE13-.LEHB13
	.uleb128 .L89-.LFB4749
	.uleb128 0
	.uleb128 .LEHB14-.LFB4749
	.uleb128 .LEHE14-.LEHB14
	.uleb128 .L90-.LFB4749
	.uleb128 0
	.uleb128 .LEHB15-.LFB4749
	.uleb128 .LEHE15-.LEHB15
	.uleb128 .L91-.LFB4749
	.uleb128 0
	.uleb128 .LEHB16-.LFB4749
	.uleb128 .LEHE16-.LEHB16
	.uleb128 .L92-.LFB4749
	.uleb128 0
	.uleb128 .LEHB17-.LFB4749
	.uleb128 .LEHE17-.LEHB17
	.uleb128 .L93-.LFB4749
	.uleb128 0
	.uleb128 .LEHB18-.LFB4749
	.uleb128 .LEHE18-.LEHB18
	.uleb128 .L75-.LFB4749
	.uleb128 0
	.uleb128 .LEHB19-.LFB4749
	.uleb128 .LEHE19-.LEHB19
	.uleb128 0
	.uleb128 0
.LLSDACSE4749:
	.section	.text.startup
	.size	_GLOBAL__sub_I__Z8copyTestPcS_yy14cudaMemcpyKind, .-_GLOBAL__sub_I__Z8copyTestPcS_yy14cudaMemcpyKind
	.section	.init_array
	.align 8
	.quad	_GLOBAL__sub_I__Z8copyTestPcS_yy14cudaMemcpyKind
	.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	_ZL10grand_size, @object
	.size	_ZL10grand_size, 152
_ZL10grand_size:
	.quad	4096
	.quad	8192
	.quad	16384
	.quad	32768
	.quad	65536
	.quad	131072
	.quad	262144
	.quad	524288
	.quad	1048576
	.quad	2097152
	.quad	4194304
	.quad	8388608
	.quad	16777216
	.quad	33554432
	.quad	67108864
	.quad	134217728
	.quad	268435456
	.quad	536870912
	.quad	1073741824
	.local	_ZL10grand_name
	.comm	_ZL10grand_name,608,32
	.section	.rodata.cst8,"aM",@progbits,8
	.align 8
.LC1:
	.long	0
	.long	1104006501
	.align 8
.LC2:
	.long	0
	.long	1083179008
	.hidden	DW.ref.__gxx_personality_v0
	.weak	DW.ref.__gxx_personality_v0
	.section	.data.rel.local.DW.ref.__gxx_personality_v0,"awG",@progbits,DW.ref.__gxx_personality_v0,comdat
	.align 8
	.type	DW.ref.__gxx_personality_v0, @object
	.size	DW.ref.__gxx_personality_v0, 8
DW.ref.__gxx_personality_v0:
	.quad	__gxx_personality_v0
	.hidden	__dso_handle
	.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 <iostream>
#include <chrono>
#include <hip/hip_runtime.h>
#include <string>
#include <iomanip>

using namespace std;

using ST = unsigned long long;

constexpr ST TOTAL_SIZE = 1 << 30; // 1 GB
constexpr ST TOTAL_SIZE_IN_BYTES = TOTAL_SIZE * sizeof(char);
constexpr ST CNT = 19;

const string grand_name[CNT] = {
	"4KB", "8KB", "16KB", "32KB", "64KB", "128KB", "256KB", "512KB", "1MB", "2MB",
	"4MB", "8MB", "16MB", "32MB", "64MB", "128MB", "256MB", "512MB", "1GB"
};

const ST grand_size[CNT] = {
	1 << 12, 1 << 13, 1 << 14, 1 << 15,
	1 << 16, 1 << 17, 1 << 18, 1 << 19,
	1 << 20, 1 << 21, 1 << 22, 1 << 23,
	1 << 24, 1 << 25, 1 << 26, 1 << 27,
	1 << 28, 1 << 29, 1 << 30
};



// granularity

double copyTest(char *dst, char *src, ST total, ST granularity, hipMemcpyKind kind) {
	ST idx = 0;
	auto t1 = chrono::system_clock::now();
	for (; idx < total; idx += granularity) {
		hipMemcpy(dst + idx, src + idx, granularity * sizeof(char), kind);
	}
	auto t2 = chrono::system_clock::now();
	chrono::duration<double> diff = t2 - t1;
	return diff.count();
}

void rangeTest(char *dst, char *src, ST total, hipMemcpyKind kind) {
	for (int i = 0; i < CNT; i++) {
		double tm = copyTest(dst, src, total, grand_size[i], kind);
		double tp = 1024 / tm; // MB/s
		cout << fixed << tp << endl;
	}
}


int main() {
	char *buf = new char[TOTAL_SIZE];
	char *buf2 = new char[TOTAL_SIZE];
	char *gpu_buf, *gpu_buf2;
	cout << "granularity: " << endl;
	for (const auto &s : grand_name)
		cout << s << endl;
	cout << endl;

	hipMalloc((void **)&gpu_buf, TOTAL_SIZE_IN_BYTES);
	hipMalloc((void **)&gpu_buf2, TOTAL_SIZE_IN_BYTES);
	memset(buf, 'a', TOTAL_SIZE);
	cout << "host to device: " << endl;
	rangeTest(gpu_buf, buf, TOTAL_SIZE, hipMemcpyHostToDevice);
	cout << endl;
	cout << "(MB/s)" << endl;

	cout << "device to device: " << endl;
	rangeTest(gpu_buf, gpu_buf2, TOTAL_SIZE, hipMemcpyDeviceToDevice);
	cout << endl;
	cout << "(MB/s)" << endl;

	cout << "device to host: " << endl;
	rangeTest(buf2, gpu_buf2, TOTAL_SIZE, hipMemcpyDeviceToHost);
	cout << endl;
	cout << buf2[32];
	cout << "(MB/s)" << endl;

	hipFree(gpu_buf);
	hipFree(gpu_buf2);
	delete[] buf;
	delete[] buf2;


//#ifdef WIN32
//	system("pause");
//#endif

	return 0;
}
	.text
	.file	"BandwidthTrail.hip"
                                        # Start of file scope inline assembly
	.globl	_ZSt21ios_base_library_initv

                                        # End of file scope inline assembly
	.section	.text.startup,"ax",@progbits
	.type	__cxx_global_array_dtor,@function # -- Begin function __cxx_global_array_dtor
__cxx_global_array_dtor:                # @__cxx_global_array_dtor
	.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	$_ZL10grand_nameB5cxx11+608, %ebx
	movq	$-608, %r14                     # imm = 0xFDA0
.LBB0_1:                                # =>This Inner Loop Header: Depth=1
	movq	-32(%rbx), %rdi
	leaq	-16(%rbx), %rax
	cmpq	%rdi, %rax
	je	.LBB0_3
# %bb.2:                                # %_ZNKSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEE11_M_is_localEv.exit.i.i
                                        #   in Loop: Header=BB0_1 Depth=1
	callq	_ZdlPv
.LBB0_3:                                # %_ZNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEED2Ev.exit
                                        #   in Loop: Header=BB0_1 Depth=1
	addq	$-32, %rbx
	addq	$32, %r14
	jne	.LBB0_1
# %bb.4:
	addq	$8, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end0:
	.size	__cxx_global_array_dtor, .Lfunc_end0-__cxx_global_array_dtor
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function _Z8copyTestPcS_yy13hipMemcpyKind
.LCPI1_0:
	.quad	0x41cdcd6500000000              # double 1.0E+9
	.text
	.globl	_Z8copyTestPcS_yy13hipMemcpyKind
	.type	_Z8copyTestPcS_yy13hipMemcpyKind,@function
_Z8copyTestPcS_yy13hipMemcpyKind:       # @_Z8copyTestPcS_yy13hipMemcpyKind
	.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
	pushq	%rax
	.cfi_def_cfa_offset 64
	.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	%r8d, %ebx
	movq	%rcx, %r14
	movq	%rdx, %r12
	movq	%rsi, %r13
	movq	%rdi, %rbp
	callq	_ZNSt6chrono3_V212system_clock3nowEv
	movq	%rax, (%rsp)                    # 8-byte Spill
	testq	%r12, %r12
	je	.LBB1_3
# %bb.1:                                # %.lr.ph.preheader
	xorl	%r15d, %r15d
.LBB1_2:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	leaq	(%r15,%rbp), %rdi
	leaq	(%r15,%r13), %rsi
	movq	%r14, %rdx
	movl	%ebx, %ecx
	callq	hipMemcpy
	addq	%r14, %r15
	cmpq	%r12, %r15
	jb	.LBB1_2
.LBB1_3:                                # %._crit_edge
	callq	_ZNSt6chrono3_V212system_clock3nowEv
	subq	(%rsp), %rax                    # 8-byte Folded Reload
	cvtsi2sd	%rax, %xmm0
	divsd	.LCPI1_0(%rip), %xmm0
	addq	$8, %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_end1:
	.size	_Z8copyTestPcS_yy13hipMemcpyKind, .Lfunc_end1-_Z8copyTestPcS_yy13hipMemcpyKind
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0                          # -- Begin function _Z9rangeTestPcS_y13hipMemcpyKind
.LCPI2_0:
	.quad	0x4090000000000000              # double 1024
	.text
	.globl	_Z9rangeTestPcS_y13hipMemcpyKind
	.type	_Z9rangeTestPcS_y13hipMemcpyKind,@function
_Z9rangeTestPcS_y13hipMemcpyKind:       # @_Z9rangeTestPcS_y13hipMemcpyKind
	.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
	pushq	%rax
	.cfi_def_cfa_offset 64
	.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	%ecx, %ebx
	movq	%rdx, %r14
	movq	%rsi, %r15
	movq	%rdi, %r12
	movq	$-152, %rbp
.LBB2_1:                                # =>This Inner Loop Header: Depth=1
	movq	_ZL10grand_size+152(%rbp), %rcx
	movq	%r12, %rdi
	movq	%r15, %rsi
	movq	%r14, %rdx
	movl	%ebx, %r8d
	callq	_Z8copyTestPcS_yy13hipMemcpyKind
	movsd	.LCPI2_0(%rip), %xmm1           # xmm1 = mem[0],zero
	divsd	%xmm0, %xmm1
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rax
	movl	_ZSt4cout+24(%rax), %ecx
	movl	$-261, %edx                     # imm = 0xFEFB
	andl	%edx, %ecx
	orl	$4, %ecx
	movl	%ecx, _ZSt4cout+24(%rax)
	movl	$_ZSt4cout, %edi
	movapd	%xmm1, %xmm0
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %r13
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%r13, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%r13, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	addq	$8, %rbp
	jne	.LBB2_1
# %bb.2:
	addq	$8, %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	_Z9rangeTestPcS_y13hipMemcpyKind, .Lfunc_end2-_Z9rangeTestPcS_y13hipMemcpyKind
	.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	$32, %rsp
	.cfi_def_cfa_offset 80
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movl	$1073741824, %edi               # imm = 0x40000000
	callq	_Znam
	movq	%rax, %rbx
	movl	$1073741824, %edi               # imm = 0x40000000
	callq	_Znam
	movq	%rax, %r14
	movl	$_ZSt4cout, %r15d
	movl	$_ZSt4cout, %edi
	movl	$.L.str.19, %esi
	movl	$13, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%r15, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movq	$-608, %r12                     # imm = 0xFDA0
.LBB3_1:                                # =>This Inner Loop Header: Depth=1
	movq	_ZL10grand_nameB5cxx11+608(%r12), %rsi
	movq	_ZL10grand_nameB5cxx11+616(%r12), %rdx
	movl	$_ZSt4cout, %edi
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	%rax, %r15
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%r15, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%r15, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	addq	$32, %r12
	jne	.LBB3_1
# %bb.2:
	movq	_ZSt4cout(%rip), %rax
	movl	$_ZSt4cout, %r13d
	movq	-24(%rax), %rdi
	addq	%r13, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	leaq	24(%rsp), %r12
	movl	$1073741824, %esi               # imm = 0x40000000
	movq	%r12, %rdi
	callq	hipMalloc
	leaq	16(%rsp), %r15
	movl	$1073741824, %esi               # imm = 0x40000000
	movq	%r15, %rdi
	callq	hipMalloc
	movl	$1073741824, %edx               # imm = 0x40000000
	movq	%rbx, %rdi
	movl	$97, %esi
	callq	memset@PLT
	movl	$_ZSt4cout, %edi
	movl	$.L.str.20, %esi
	movl	$16, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%r13, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movq	(%r12), %rdi
	movl	$1073741824, %edx               # imm = 0x40000000
	movq	%rbx, %rsi
	movl	$1, %ecx
	callq	_Z9rangeTestPcS_y13hipMemcpyKind
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%r13, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$_ZSt4cout, %edi
	movl	$.L.str.21, %esi
	movl	$6, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%r13, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$_ZSt4cout, %edi
	movl	$.L.str.22, %esi
	movl	$18, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%r13, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movq	(%r12), %rdi
	movq	(%r15), %rsi
	movl	$1073741824, %edx               # imm = 0x40000000
	movl	$3, %ecx
	callq	_Z9rangeTestPcS_y13hipMemcpyKind
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%r13, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$_ZSt4cout, %edi
	movl	$.L.str.21, %esi
	movl	$6, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%r13, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movl	$_ZSt4cout, %edi
	movl	$.L.str.23, %esi
	movl	$16, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%r13, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movq	(%r15), %rsi
	movl	$1073741824, %edx               # imm = 0x40000000
	movq	%r14, %rdi
	movl	$2, %ecx
	callq	_Z9rangeTestPcS_y13hipMemcpyKind
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%r13, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movb	32(%r14), %al
	movb	%al, 15(%rsp)
	movq	_ZSt4cout(%rip), %rcx
	movq	-24(%rcx), %rcx
	cmpq	$0, _ZSt4cout+16(%rcx)
	je	.LBB3_4
# %bb.3:
	leaq	15(%rsp), %rsi
	movl	$_ZSt4cout, %edi
	movl	$1, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	jmp	.LBB3_5
.LBB3_4:
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
.LBB3_5:                                # %_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_c.exit
	movl	$_ZSt4cout, %r15d
	movl	$_ZSt4cout, %edi
	movl	$.L.str.21, %esi
	movl	$6, %edx
	callq	_ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
	movq	_ZSt4cout(%rip), %rax
	movq	-24(%rax), %rdi
	addq	%r15, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movl	$_ZSt4cout, %edi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movq	24(%rsp), %rdi
	callq	hipFree
	movq	16(%rsp), %rdi
	callq	hipFree
	movq	%rbx, %rdi
	callq	_ZdaPv
	movq	%r14, %rdi
	callq	_ZdaPv
	xorl	%eax, %eax
	addq	$32, %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_end3:
	.size	main, .Lfunc_end3-main
	.cfi_endproc
                                        # -- End function
	.section	.text.startup,"ax",@progbits
	.type	_GLOBAL__sub_I_BandwidthTrail.hip,@function # -- Begin function _GLOBAL__sub_I_BandwidthTrail.hip
_GLOBAL__sub_I_BandwidthTrail.hip:      # @_GLOBAL__sub_I_BandwidthTrail.hip
	.cfi_startproc
# %bb.0:                                # %__cxx_global_var_init.exit
	movq	$_ZL10grand_nameB5cxx11+16, _ZL10grand_nameB5cxx11(%rip)
	movl	$4344628, _ZL10grand_nameB5cxx11+16(%rip) # imm = 0x424B34
	movl	$3, %eax
	movq	%rax, _ZL10grand_nameB5cxx11+8(%rip)
	movq	$_ZL10grand_nameB5cxx11+48, _ZL10grand_nameB5cxx11+32(%rip)
	movl	$4344632, _ZL10grand_nameB5cxx11+48(%rip) # imm = 0x424B38
	movq	%rax, _ZL10grand_nameB5cxx11+40(%rip)
	movq	$_ZL10grand_nameB5cxx11+80, _ZL10grand_nameB5cxx11+64(%rip)
	movl	$1112225329, _ZL10grand_nameB5cxx11+80(%rip) # imm = 0x424B3631
	movl	$4, %edx
	movq	%rdx, _ZL10grand_nameB5cxx11+72(%rip)
	xorl	%ecx, %ecx
	movb	%cl, _ZL10grand_nameB5cxx11+84(%rip)
	movq	$_ZL10grand_nameB5cxx11+112, _ZL10grand_nameB5cxx11+96(%rip)
	movl	$1112224307, _ZL10grand_nameB5cxx11+112(%rip) # imm = 0x424B3233
	movq	%rdx, _ZL10grand_nameB5cxx11+104(%rip)
	movb	%cl, _ZL10grand_nameB5cxx11+116(%rip)
	movq	$_ZL10grand_nameB5cxx11+144, _ZL10grand_nameB5cxx11+128(%rip)
	movl	$1112224822, _ZL10grand_nameB5cxx11+144(%rip) # imm = 0x424B3436
	movq	%rdx, _ZL10grand_nameB5cxx11+136(%rip)
	movb	%cl, _ZL10grand_nameB5cxx11+148(%rip)
	movq	$_ZL10grand_nameB5cxx11+176, _ZL10grand_nameB5cxx11+160(%rip)
	movw	$66, %di
	movw	%di, _ZL10grand_nameB5cxx11+180(%rip)
	movl	$1261974065, _ZL10grand_nameB5cxx11+176(%rip) # imm = 0x4B383231
	movl	$5, %esi
	movq	%rsi, _ZL10grand_nameB5cxx11+168(%rip)
	movq	$_ZL10grand_nameB5cxx11+208, _ZL10grand_nameB5cxx11+192(%rip)
	movw	%di, _ZL10grand_nameB5cxx11+212(%rip)
	movl	$1261843762, _ZL10grand_nameB5cxx11+208(%rip) # imm = 0x4B363532
	movq	%rsi, _ZL10grand_nameB5cxx11+200(%rip)
	movq	$_ZL10grand_nameB5cxx11+240, _ZL10grand_nameB5cxx11+224(%rip)
	movw	%di, _ZL10grand_nameB5cxx11+244(%rip)
	movl	$1261580597, _ZL10grand_nameB5cxx11+240(%rip) # imm = 0x4B323135
	movq	%rsi, _ZL10grand_nameB5cxx11+232(%rip)
	movq	$_ZL10grand_nameB5cxx11+272, _ZL10grand_nameB5cxx11+256(%rip)
	movl	$4345137, _ZL10grand_nameB5cxx11+272(%rip) # imm = 0x424D31
	movq	%rax, _ZL10grand_nameB5cxx11+264(%rip)
	movq	$_ZL10grand_nameB5cxx11+304, _ZL10grand_nameB5cxx11+288(%rip)
	movl	$4345138, _ZL10grand_nameB5cxx11+304(%rip) # imm = 0x424D32
	movq	%rax, _ZL10grand_nameB5cxx11+296(%rip)
	movq	$_ZL10grand_nameB5cxx11+336, _ZL10grand_nameB5cxx11+320(%rip)
	movb	$66, %dil
	movb	%dil, _ZL10grand_nameB5cxx11+338(%rip)
	movw	$19764, _ZL10grand_nameB5cxx11+336(%rip) # imm = 0x4D34
	movq	%rax, _ZL10grand_nameB5cxx11+328(%rip)
	movb	%cl, _ZL10grand_nameB5cxx11+339(%rip)
	movq	$_ZL10grand_nameB5cxx11+368, _ZL10grand_nameB5cxx11+352(%rip)
	movb	%dil, _ZL10grand_nameB5cxx11+370(%rip)
	movw	$19768, _ZL10grand_nameB5cxx11+368(%rip) # imm = 0x4D38
	movq	%rax, _ZL10grand_nameB5cxx11+360(%rip)
	movb	%cl, _ZL10grand_nameB5cxx11+371(%rip)
	movq	$_ZL10grand_nameB5cxx11+400, _ZL10grand_nameB5cxx11+384(%rip)
	movl	$1112356401, _ZL10grand_nameB5cxx11+400(%rip) # imm = 0x424D3631
	movq	%rdx, _ZL10grand_nameB5cxx11+392(%rip)
	movb	%cl, _ZL10grand_nameB5cxx11+404(%rip)
	movq	$_ZL10grand_nameB5cxx11+432, _ZL10grand_nameB5cxx11+416(%rip)
	movl	$1112355379, _ZL10grand_nameB5cxx11+432(%rip) # imm = 0x424D3233
	movq	%rdx, _ZL10grand_nameB5cxx11+424(%rip)
	movb	%cl, _ZL10grand_nameB5cxx11+436(%rip)
	movq	$_ZL10grand_nameB5cxx11+464, _ZL10grand_nameB5cxx11+448(%rip)
	movl	$1112355894, _ZL10grand_nameB5cxx11+464(%rip) # imm = 0x424D3436
	movq	%rdx, _ZL10grand_nameB5cxx11+456(%rip)
	movb	%cl, _ZL10grand_nameB5cxx11+468(%rip)
	movq	$_ZL10grand_nameB5cxx11+496, _ZL10grand_nameB5cxx11+480(%rip)
	movb	%dil, _ZL10grand_nameB5cxx11+500(%rip)
	movl	$1295528497, _ZL10grand_nameB5cxx11+496(%rip) # imm = 0x4D383231
	movq	%rsi, _ZL10grand_nameB5cxx11+488(%rip)
	movb	%cl, _ZL10grand_nameB5cxx11+501(%rip)
	movq	$_ZL10grand_nameB5cxx11+528, _ZL10grand_nameB5cxx11+512(%rip)
	movb	%dil, _ZL10grand_nameB5cxx11+532(%rip)
	movl	$1295398194, _ZL10grand_nameB5cxx11+528(%rip) # imm = 0x4D363532
	movq	%rsi, _ZL10grand_nameB5cxx11+520(%rip)
	movb	%cl, _ZL10grand_nameB5cxx11+533(%rip)
	movq	$_ZL10grand_nameB5cxx11+560, _ZL10grand_nameB5cxx11+544(%rip)
	movb	%dil, _ZL10grand_nameB5cxx11+564(%rip)
	movl	$1295135029, _ZL10grand_nameB5cxx11+560(%rip) # imm = 0x4D323135
	movq	%rsi, _ZL10grand_nameB5cxx11+552(%rip)
	movb	%cl, _ZL10grand_nameB5cxx11+565(%rip)
	movq	$_ZL10grand_nameB5cxx11+592, _ZL10grand_nameB5cxx11+576(%rip)
	movb	%dil, _ZL10grand_nameB5cxx11+594(%rip)
	movw	$18225, _ZL10grand_nameB5cxx11+592(%rip) # imm = 0x4731
	movq	%rax, _ZL10grand_nameB5cxx11+584(%rip)
	movb	%cl, _ZL10grand_nameB5cxx11+595(%rip)
	movl	$__cxx_global_array_dtor, %edi
	movl	$__dso_handle, %edx
	xorl	%esi, %esi
	jmp	__cxa_atexit                    # TAILCALL
.Lfunc_end4:
	.size	_GLOBAL__sub_I_BandwidthTrail.hip, .Lfunc_end4-_GLOBAL__sub_I_BandwidthTrail.hip
	.cfi_endproc
                                        # -- End function
	.type	_ZL10grand_nameB5cxx11,@object  # @_ZL10grand_nameB5cxx11
	.local	_ZL10grand_nameB5cxx11
	.comm	_ZL10grand_nameB5cxx11,608,16
	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"4KB"
	.size	.L.str, 4

	.type	.L.str.1,@object                # @.str.1
.L.str.1:
	.asciz	"8KB"
	.size	.L.str.1, 4

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

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

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

	.type	.L.str.8,@object                # @.str.8
.L.str.8:
	.asciz	"1MB"
	.size	.L.str.8, 4

	.type	.L.str.9,@object                # @.str.9
.L.str.9:
	.asciz	"2MB"
	.size	.L.str.9, 4

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

	.type	.L.str.11,@object               # @.str.11
.L.str.11:
	.asciz	"8MB"
	.size	.L.str.11, 4

	.type	.L.str.15,@object               # @.str.15
.L.str.15:
	.asciz	"128MB"
	.size	.L.str.15, 6

	.type	.L.str.16,@object               # @.str.16
.L.str.16:
	.asciz	"256MB"
	.size	.L.str.16, 6

	.type	.L.str.17,@object               # @.str.17
.L.str.17:
	.asciz	"512MB"
	.size	.L.str.17, 6

	.type	.L.str.18,@object               # @.str.18
.L.str.18:
	.asciz	"1GB"
	.size	.L.str.18, 4

	.hidden	__dso_handle
	.type	_ZL10grand_size,@object         # @_ZL10grand_size
	.section	.rodata,"a",@progbits
	.p2align	4, 0x0
_ZL10grand_size:
	.quad	4096                            # 0x1000
	.quad	8192                            # 0x2000
	.quad	16384                           # 0x4000
	.quad	32768                           # 0x8000
	.quad	65536                           # 0x10000
	.quad	131072                          # 0x20000
	.quad	262144                          # 0x40000
	.quad	524288                          # 0x80000
	.quad	1048576                         # 0x100000
	.quad	2097152                         # 0x200000
	.quad	4194304                         # 0x400000
	.quad	8388608                         # 0x800000
	.quad	16777216                        # 0x1000000
	.quad	33554432                        # 0x2000000
	.quad	67108864                        # 0x4000000
	.quad	134217728                       # 0x8000000
	.quad	268435456                       # 0x10000000
	.quad	536870912                       # 0x20000000
	.quad	1073741824                      # 0x40000000
	.size	_ZL10grand_size, 152

	.type	.L.str.19,@object               # @.str.19
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str.19:
	.asciz	"granularity: "
	.size	.L.str.19, 14

	.type	.L.str.20,@object               # @.str.20
.L.str.20:
	.asciz	"host to device: "
	.size	.L.str.20, 17

	.type	.L.str.21,@object               # @.str.21
.L.str.21:
	.asciz	"(MB/s)"
	.size	.L.str.21, 7

	.type	.L.str.22,@object               # @.str.22
.L.str.22:
	.asciz	"device to device: "
	.size	.L.str.22, 19

	.type	.L.str.23,@object               # @.str.23
.L.str.23:
	.asciz	"device to host: "
	.size	.L.str.23, 17

	.section	.init_array,"aw",@init_array
	.p2align	3, 0x0
	.quad	_GLOBAL__sub_I_BandwidthTrail.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 __gxx_personality_v0
	.addrsig_sym __cxx_global_array_dtor
	.addrsig_sym _GLOBAL__sub_I_BandwidthTrail.hip
	.addrsig_sym _ZL10grand_nameB5cxx11
	.addrsig_sym __dso_handle
	.addrsig_sym _ZSt4cout
	.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
