/*
compile : nvcc malloc_test.cu -o malloc_test
exec : ./malloc_test 512 1
starting 512MB , increase 1MB 
if you have multiple GPU, use below
export CUDA_VISIBLE_DEVICES=3;./malloc_test 512 1

MB convert table would be : 
GB	MB
1	1024
2	2048
3	3072
4	4096
5	5120
6	6144
7	7168
8	8192
9	9216
10	10240
11	11264
12	12288
13	13312
14	14336
15	15360
16	16384
17	17408
18	18432
19	19456
20	20480
21	21504
22	22528
23	23552
24	24576
25	25600
26	26624
27	27648
28	28672
29	29696
30	30720
31	31744
32	32768
*/

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

int main(int argc, char* argv[]) {

	int *devPtr= NULL;
	size_t mem_size, free, total , start, inc, size ;
	float free_f, total_f, used_f, used_start_f, used_now_f;

        cudaMemGetInfo(&free,&total);
        free_f  = float(free)   / (1024*1024) ; 
        total_f = float(total)  / (1024*1024) ;
        used_f =  total_f-free_f ;
	used_start_f  = used_f;

	if(argc<3){
	printf(" plz use below command \n ./malloc_test 1024 10  \n to  malloc 1024MB and increment would be 10MB \n");
	printf("\n11GB 11264MB\n15GB 15360MB \n23GB 23552MB\n31GB 31744MB\n");
	printf("\ncurrent free memory is %.1f\n", free_f);

	return 0;
	}

        start =  atoi(argv[1]) / sizeof(int) ;
        inc = atoi(argv[2])  ;
        size = start;
	
        printf("------------------------------------------------------------------\n");
        printf("\tTotal(MB)=\tFree(MB)+\tinit(MB)+\tAlloc(MB)\n"); 
        printf("0\t%.1f    =\t%.1f+    \t%.1f     \t <------  initial used memory \n", total_f,free_f, used_f );
        printf("------------------------------------------------------------------\n");
	int i = 0;

	do {
		
		mem_size = sizeof(int) * size * (1024*1024) + (inc*i) * (1024*1024) ; 
		cudaMalloc(&devPtr, mem_size ); // MB
                cudaMemGetInfo(&free,&total);
	        free_f  = float(free)   / (1024*1024) ;
        	total_f = float(total)  / (1024*1024) ;
	        used_f =  total_f-free_f ;
		used_now_f = (float)mem_size/(1024*1024)   ; 
        printf("%d\t%.1f    =\t%.1f+    \t%.1f+    \t%.1f \n", i, total_f, free_f, used_start_f, used_now_f);
		if(devPtr == NULL) {
			printf("couldn't allocate %.1f MB ", used_now_f);
			printf("Err : %s\n", cudaGetErrorString(cudaGetLastError()) );
			return 0;	
		} else {
			//printf("Allocated %d int's.\n", int(size));
		}
		cudaFree(devPtr);
		size = (size* sizeof(int) + inc )/sizeof(int) ;  
		mem_size = sizeof(int) * size  ;
		
		i=i+1;
	} while(1);

}

	.file	"tmpxft_00384546_00000000-6_malloc_test.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2030:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2030:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC1:
	.string	" plz use below command \n ./malloc_test 1024 10  \n to  malloc 1024MB and increment would be 10MB \n"
.LC2:
	.string	"\n11GB 11264MB\n15GB 15360MB \n23GB 23552MB\n31GB 31744MB\n"
.LC3:
	.string	"\ncurrent free memory is %.1f\n"
.LC4:
	.string	"------------------------------------------------------------------\n"
.LC5:
	.string	"\tTotal(MB)=\tFree(MB)+\tinit(MB)+\tAlloc(MB)\n"
.LC6:
	.string	"0\t%.1f    =\t%.1f+    \t%.1f     \t <------  initial used memory \n"
.LC7:
	.string	"%d\t%.1f    =\t%.1f+    \t%.1f+    \t%.1f \n"
.LC8:
	.string	"couldn't allocate %.1f MB "
.LC9:
	.string	"Err : %s\n"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2027:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movq	%rsi, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movl	%edi, %ebx
	subq	$72, %rsp
	.cfi_def_cfa_offset 128
	movq	%fs:40, %rax
	movq	%rax, 56(%rsp)
	xorl	%eax, %eax
	xorl	%eax, %eax
	leaq	40(%rsp), %r14
	movq	%rax, 32(%rsp)
	leaq	48(%rsp), %rax
	movq	%r14, %rdi
	movq	%rax, %rsi
	movq	%rax, 8(%rsp)
	call	cudaMemGetInfo@PLT
	movq	40(%rsp), %rax
	testq	%rax, %rax
	js	.L3
	cvtsi2ssq	%rax, %xmm2
	jmp	.L4
.L3:
	movq	%rax, %rdx
	andl	$1, %eax
	shrq	%rdx
	orq	%rax, %rdx
	cvtsi2ssq	%rdx, %xmm2
	addss	%xmm2, %xmm2
.L4:
	mulss	.LC0(%rip), %xmm2
	cmpl	$2, %ebx
	movq	48(%rsp), %rax
	cvtss2sd	%xmm2, %xmm1
	jg	.L5
	leaq	.LC1(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	movsd	%xmm1, (%rsp)
	call	__printf_chk@PLT
	leaq	.LC2(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movsd	(%rsp), %xmm1
	leaq	.LC3(%rip), %rsi
	movb	$1, %al
	movl	$2, %edi
	movaps	%xmm1, %xmm0
	call	__printf_chk@PLT
	jmp	.L6
.L5:
	testq	%rax, %rax
	js	.L7
	cvtsi2ssq	%rax, %xmm0
	jmp	.L8
.L7:
	movq	%rax, %rdx
	andl	$1, %eax
	shrq	%rdx
	orq	%rax, %rdx
	cvtsi2ssq	%rdx, %xmm0
	addss	%xmm0, %xmm0
.L8:
	mulss	.LC0(%rip), %xmm0
	movq	8(%rbp), %rdi
	xorl	%r13d, %r13d
	movsd	%xmm1, 24(%rsp)
	movss	%xmm2, 16(%rsp)
	leaq	.LC4(%rip), %r12
	leaq	.LC7(%rip), %r15
	movss	%xmm0, (%rsp)
	call	atoi@PLT
	movq	16(%rbp), %rdi
	movslq	%eax, %rbx
	call	atoi@PLT
	movq	%r12, %rsi
	movl	$2, %edi
	shrq	$2, %rbx
	movslq	%eax, %rbp
	xorl	%eax, %eax
	call	__printf_chk@PLT
	leaq	.LC5(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movss	(%rsp), %xmm0
	movb	$3, %al
	movss	16(%rsp), %xmm2
	movsd	24(%rsp), %xmm1
	leaq	.LC6(%rip), %rsi
	movl	$2, %edi
	movaps	%xmm0, %xmm3
	cvtss2sd	%xmm0, %xmm0
	subss	%xmm2, %xmm3
	cvtss2sd	%xmm3, %xmm4
	movaps	%xmm4, %xmm2
	movsd	%xmm4, (%rsp)
	call	__printf_chk@PLT
	movq	%r12, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	xorl	%r12d, %r12d
.L16:
	salq	$2, %rbx
	leaq	32(%rsp), %rdi
	leaq	(%rbx,%r13), %rdx
	salq	$20, %rdx
	movq	%rdx, %rsi
	movq	%rdx, 16(%rsp)
	call	cudaMalloc@PLT
	movq	8(%rsp), %rsi
	movq	%r14, %rdi
	call	cudaMemGetInfo@PLT
	movq	16(%rsp), %rdx
	testq	%rdx, %rdx
	js	.L9
	cvtsi2ssq	%rdx, %xmm3
	jmp	.L10
.L9:
	shrq	%rdx
	cvtsi2ssq	%rdx, %xmm3
	addss	%xmm3, %xmm3
.L10:
	mulss	.LC0(%rip), %xmm3
	movq	40(%rsp), %rax
	cvtss2sd	%xmm3, %xmm3
	testq	%rax, %rax
	js	.L11
	cvtsi2ssq	%rax, %xmm1
	jmp	.L12
.L11:
	movq	%rax, %rdx
	andl	$1, %eax
	shrq	%rdx
	orq	%rax, %rdx
	cvtsi2ssq	%rdx, %xmm1
	addss	%xmm1, %xmm1
.L12:
	mulss	.LC0(%rip), %xmm1
	movq	48(%rsp), %rax
	cvtss2sd	%xmm1, %xmm1
	testq	%rax, %rax
	js	.L13
	cvtsi2ssq	%rax, %xmm0
	jmp	.L14
.L13:
	movq	%rax, %rdx
	andl	$1, %eax
	shrq	%rdx
	orq	%rax, %rdx
	cvtsi2ssq	%rdx, %xmm0
	addss	%xmm0, %xmm0
.L14:
	mulss	.LC0(%rip), %xmm0
	movl	%r12d, %edx
	movq	%r15, %rsi
	movb	$4, %al
	movsd	(%rsp), %xmm2
	movl	$2, %edi
	movsd	%xmm3, 16(%rsp)
	addq	%rbp, %r13
	cvtss2sd	%xmm0, %xmm0
	call	__printf_chk@PLT
	movq	32(%rsp), %rdi
	movsd	16(%rsp), %xmm3
	testq	%rdi, %rdi
	jne	.L15
	movaps	%xmm3, %xmm0
	leaq	.LC8(%rip), %rsi
	movl	$2, %edi
	movb	$1, %al
	call	__printf_chk@PLT
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	call	cudaGetErrorString@PLT
	leaq	.LC9(%rip), %rsi
	movl	$2, %edi
	movq	%rax, %rdx
	xorl	%eax, %eax
	call	__printf_chk@PLT
	jmp	.L6
.L15:
	call	cudaFree@PLT
	addq	%rbp, %rbx
	incl	%r12d
	shrq	$2, %rbx
	jmp	.L16
.L6:
	movq	56(%rsp), %rax
	subq	%fs:40, %rax
	je	.L17
	call	__stack_chk_fail@PLT
.L17:
	addq	$72, %rsp
	.cfi_def_cfa_offset 56
	xorl	%eax, %eax
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2027:
	.size	main, .-main
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2053:
	.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
.LFE2053:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.nvFatBinSegment,"aw"
	.align 8
	.type	_ZL15__fatDeviceText, @object
	.size	_ZL15__fatDeviceText, 24
_ZL15__fatDeviceText:
	.long	1180844977
	.long	1
	.quad	fatbinData
	.quad	0
	.local	_ZL20__cudaFatCubinHandle
	.comm	_ZL20__cudaFatCubinHandle,8,8
	.section	.rodata.cst4,"aM",@progbits,4
	.align 4
.LC0:
	.long	897581056
	.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:

/*
compile : nvcc malloc_test.cu -o malloc_test
exec : ./malloc_test 512 1
starting 512MB , increase 1MB 
if you have multiple GPU, use below
export CUDA_VISIBLE_DEVICES=3;./malloc_test 512 1

MB convert table would be : 
GB	MB
1	1024
2	2048
3	3072
4	4096
5	5120
6	6144
7	7168
8	8192
9	9216
10	10240
11	11264
12	12288
13	13312
14	14336
15	15360
16	16384
17	17408
18	18432
19	19456
20	20480
21	21504
22	22528
23	23552
24	24576
25	25600
26	26624
27	27648
28	28672
29	29696
30	30720
31	31744
32	32768
*/


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

int main(int argc, char* argv[]) {

	int *devPtr= NULL;
	size_t mem_size, free, total , start, inc, size ;
	float free_f, total_f, used_f, used_start_f, used_now_f;

        hipMemGetInfo(&free,&total);
        free_f  = float(free)   / (1024*1024) ; 
        total_f = float(total)  / (1024*1024) ;
        used_f =  total_f-free_f ;
	used_start_f  = used_f;

	if(argc<3){
	printf(" plz use below command \n ./malloc_test 1024 10  \n to  malloc 1024MB and increment would be 10MB \n");
	printf("\n11GB 11264MB\n15GB 15360MB \n23GB 23552MB\n31GB 31744MB\n");
	printf("\ncurrent free memory is %.1f\n", free_f);

	return 0;
	}

        start =  atoi(argv[1]) / sizeof(int) ;
        inc = atoi(argv[2])  ;
        size = start;
	
        printf("------------------------------------------------------------------\n");
        printf("\tTotal(MB)=\tFree(MB)+\tinit(MB)+\tAlloc(MB)\n"); 
        printf("0\t%.1f    =\t%.1f+    \t%.1f     \t <------  initial used memory \n", total_f,free_f, used_f );
        printf("------------------------------------------------------------------\n");
	int i = 0;

	do {
		
		mem_size = sizeof(int) * size * (1024*1024) + (inc*i) * (1024*1024) ; 
		hipMalloc(&devPtr, mem_size ); // MB
                hipMemGetInfo(&free,&total);
	        free_f  = float(free)   / (1024*1024) ;
        	total_f = float(total)  / (1024*1024) ;
	        used_f =  total_f-free_f ;
		used_now_f = (float)mem_size/(1024*1024)   ; 
        printf("%d\t%.1f    =\t%.1f+    \t%.1f+    \t%.1f \n", i, total_f, free_f, used_start_f, used_now_f);
		if(devPtr == NULL) {
			printf("couldn't allocate %.1f MB ", used_now_f);
			printf("Err : %s\n", hipGetErrorString(hipGetLastError()) );
			return 0;	
		} else {
			//printf("Allocated %d int's.\n", int(size));
		}
		hipFree(devPtr);
		size = (size* sizeof(int) + inc )/sizeof(int) ;  
		mem_size = sizeof(int) * size  ;
		
		i=i+1;
	} while(1);

}

	.text
	.file	"malloc_test.hip"
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0                          # -- Begin function main
.LCPI0_0:
	.long	0x35800000                      # float 9.53674316E-7
	.text
	.globl	main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%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, %rbx
	movl	%edi, %ebp
	movq	$0, 48(%rsp)
	leaq	40(%rsp), %r14
	leaq	24(%rsp), %rsi
	movq	%r14, %rdi
	callq	hipMemGetInfo
	movq	(%r14), %rax
	testq	%rax, %rax
	js	.LBB0_1
# %bb.2:
	cvtsi2ss	%rax, %xmm0
	jmp	.LBB0_3
.LBB0_1:
	movq	%rax, %rcx
	shrq	%rcx
	andl	$1, %eax
	orq	%rcx, %rax
	cvtsi2ss	%rax, %xmm0
	addss	%xmm0, %xmm0
.LBB0_3:
	mulss	.LCPI0_0(%rip), %xmm0
	cmpl	$2, %ebp
	movss	%xmm0, 8(%rsp)                  # 4-byte Spill
	jg	.LBB0_5
# %bb.4:
	movl	$.Lstr.3, %edi
	callq	puts@PLT
	movl	$.Lstr.4, %edi
	callq	puts@PLT
	xorps	%xmm0, %xmm0
	cvtss2sd	8(%rsp), %xmm0                  # 4-byte Folded Reload
	movl	$.L.str.2, %edi
	movb	$1, %al
	callq	printf
	jmp	.LBB0_20
.LBB0_5:
	movq	24(%rsp), %rax
	testq	%rax, %rax
	js	.LBB0_6
# %bb.7:
	cvtsi2ss	%rax, %xmm1
	jmp	.LBB0_8
.LBB0_6:
	movq	%rax, %rcx
	shrq	%rcx
	andl	$1, %eax
	orq	%rcx, %rax
	cvtsi2ss	%rax, %xmm1
	addss	%xmm1, %xmm1
.LBB0_8:
	mulss	.LCPI0_0(%rip), %xmm1
	movss	%xmm1, 36(%rsp)                 # 4-byte Spill
	movaps	%xmm1, %xmm2
	subss	%xmm0, %xmm2
	movss	%xmm2, 16(%rsp)                 # 4-byte Spill
	movq	8(%rbx), %rdi
	callq	atoi
	movslq	%eax, %rbp
	shrq	$2, %rbp
	movq	16(%rbx), %rdi
	callq	atoi
	movslq	%eax, %r15
	movl	$.Lstr.2, %edi
	callq	puts@PLT
	movl	$.Lstr.1, %edi
	callq	puts@PLT
	xorps	%xmm0, %xmm0
	cvtss2sd	36(%rsp), %xmm0                 # 4-byte Folded Reload
	xorps	%xmm1, %xmm1
	cvtss2sd	8(%rsp), %xmm1                  # 4-byte Folded Reload
	xorps	%xmm2, %xmm2
	cvtss2sd	16(%rsp), %xmm2                 # 4-byte Folded Reload
	movl	$.L.str.5, %edi
	movsd	%xmm2, 8(%rsp)                  # 8-byte Spill
	movb	$3, %al
	callq	printf
	movl	$.Lstr.2, %edi
	callq	puts@PLT
	movq	%r15, %r14
	shlq	$20, %r14
	shrq	$2, %r15
	xorl	%ebx, %ebx
	xorl	%r12d, %r12d
.LBB0_9:                                # =>This Inner Loop Header: Depth=1
	movq	%rbp, %r13
	shlq	$22, %r13
	addq	%r12, %r13
	leaq	48(%rsp), %rdi
	movq	%r13, %rsi
	callq	hipMalloc
	leaq	40(%rsp), %rdi
	leaq	24(%rsp), %rsi
	callq	hipMemGetInfo
	movq	40(%rsp), %rax
	testq	%rax, %rax
	js	.LBB0_10
# %bb.11:                               #   in Loop: Header=BB0_9 Depth=1
	xorps	%xmm1, %xmm1
	cvtsi2ss	%rax, %xmm1
	jmp	.LBB0_12
.LBB0_10:                               #   in Loop: Header=BB0_9 Depth=1
	movq	%rax, %rcx
	shrq	%rcx
	andl	$1, %eax
	orq	%rcx, %rax
	xorps	%xmm1, %xmm1
	cvtsi2ss	%rax, %xmm1
	addss	%xmm1, %xmm1
.LBB0_12:                               #   in Loop: Header=BB0_9 Depth=1
	movq	24(%rsp), %rax
	testq	%rax, %rax
	js	.LBB0_13
# %bb.14:                               #   in Loop: Header=BB0_9 Depth=1
	xorps	%xmm0, %xmm0
	cvtsi2ss	%rax, %xmm0
	jmp	.LBB0_15
.LBB0_13:                               #   in Loop: Header=BB0_9 Depth=1
	movq	%rax, %rcx
	shrq	%rcx
	andl	$1, %eax
	orq	%rcx, %rax
	xorps	%xmm0, %xmm0
	cvtsi2ss	%rax, %xmm0
	addss	%xmm0, %xmm0
.LBB0_15:                               #   in Loop: Header=BB0_9 Depth=1
	testq	%r13, %r13
	js	.LBB0_16
# %bb.17:                               #   in Loop: Header=BB0_9 Depth=1
	xorps	%xmm2, %xmm2
	cvtsi2ss	%r13, %xmm2
	jmp	.LBB0_18
.LBB0_16:                               #   in Loop: Header=BB0_9 Depth=1
	movq	%r13, %rax
	shrq	%rax
	andl	$1, %r13d
	orq	%rax, %r13
	xorps	%xmm2, %xmm2
	cvtsi2ss	%r13, %xmm2
	addss	%xmm2, %xmm2
.LBB0_18:                               #   in Loop: Header=BB0_9 Depth=1
	movss	.LCPI0_0(%rip), %xmm3           # xmm3 = mem[0],zero,zero,zero
	mulss	%xmm3, %xmm1
	mulss	%xmm3, %xmm0
	mulss	%xmm3, %xmm2
	cvtss2sd	%xmm0, %xmm0
	cvtss2sd	%xmm1, %xmm1
	xorps	%xmm3, %xmm3
	cvtss2sd	%xmm2, %xmm3
	movl	$.L.str.6, %edi
	movl	%ebx, %esi
	movsd	8(%rsp), %xmm2                  # 8-byte Reload
                                        # xmm2 = mem[0],zero
	movsd	%xmm3, 16(%rsp)                 # 8-byte Spill
	movb	$4, %al
	callq	printf
	movq	48(%rsp), %rdi
	testq	%rdi, %rdi
	je	.LBB0_19
# %bb.21:                               #   in Loop: Header=BB0_9 Depth=1
	callq	hipFree
	addq	%r15, %rbp
	movabsq	$4611686018427387903, %rax      # imm = 0x3FFFFFFFFFFFFFFF
	andq	%rax, %rbp
	addq	%r14, %r12
	incl	%ebx
	jmp	.LBB0_9
.LBB0_19:
	movl	$.L.str.7, %edi
	movsd	16(%rsp), %xmm0                 # 8-byte Reload
                                        # xmm0 = mem[0],zero
	movb	$1, %al
	callq	printf
	callq	hipGetLastError
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.8, %edi
	movq	%rax, %rsi
	xorl	%eax, %eax
	callq	printf
.LBB0_20:
	xorl	%eax, %eax
	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_end0:
	.size	main, .Lfunc_end0-main
	.cfi_endproc
                                        # -- End function
	.type	.L.str.2,@object                # @.str.2
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str.2:
	.asciz	"\ncurrent free memory is %.1f\n"
	.size	.L.str.2, 30

	.type	.L.str.5,@object                # @.str.5
.L.str.5:
	.asciz	"0\t%.1f    =\t%.1f+    \t%.1f     \t <------  initial used memory \n"
	.size	.L.str.5, 64

	.type	.L.str.6,@object                # @.str.6
.L.str.6:
	.asciz	"%d\t%.1f    =\t%.1f+    \t%.1f+    \t%.1f \n"
	.size	.L.str.6, 40

	.type	.L.str.7,@object                # @.str.7
.L.str.7:
	.asciz	"couldn't allocate %.1f MB "
	.size	.L.str.7, 27

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

	.type	__hip_cuid_,@object             # @__hip_cuid_
	.bss
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               # 0x0
	.size	__hip_cuid_, 1

	.type	.Lstr.1,@object                 # @str.1
	.section	.rodata.str1.1,"aMS",@progbits,1
.Lstr.1:
	.asciz	"\tTotal(MB)=\tFree(MB)+\tinit(MB)+\tAlloc(MB)"
	.size	.Lstr.1, 42

	.type	.Lstr.2,@object                 # @str.2
.Lstr.2:
	.asciz	"------------------------------------------------------------------"
	.size	.Lstr.2, 67

	.type	.Lstr.3,@object                 # @str.3
.Lstr.3:
	.asciz	" plz use below command \n ./malloc_test 1024 10  \n to  malloc 1024MB and increment would be 10MB "
	.size	.Lstr.3, 97

	.type	.Lstr.4,@object                 # @str.4
.Lstr.4:
	.asciz	"\n11GB 11264MB\n15GB 15360MB \n23GB 23552MB\n31GB 31744MB"
	.size	.Lstr.4, 54

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

	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.amdgpu_metadata
---
amdhsa.kernels:  []
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
