/*
 * bandwidth-test, used for plotting
 * can use pageable or pinned memory
 */
#include <stdio.h>
#include <getopt.h>
#include <cuda_runtime.h>

#define YELLOW "\e[1;33m"
#define RED    "\e[1;31m"
#define WHITE  "\e[1;00m"



typedef struct {
    float mseconds_h2d;
    float mseconds_d2h;
    float bandwidth_h2d;
    float bandwidth_d2h;
    bool err;
} Profiling;

typedef enum {
    PLOT,
    HUMAN
} Format;

typedef enum {
    PINNED,
    PAGEABLE
} Mode;



struct option long_options[] = {
    {"num"      , required_argument, 0, 0},
    {"mode"     , required_argument, 0, 0},
    {"format"   , required_argument, 0, 0},
    {"help"     , 0                , 0, 0},
    {0          , 0                , 0, 0}
};

float calcBandwidth(size_t size, float mseconds) {
    return (size/((float)1024*1024*1024)) / (mseconds / 1000.0f);
}

/*! perform a profiled memory transfer
 *
 *  with pinned or pageable memory using size bytes
 *
 */
Profiling profileTransfer(size_t size, Mode mode, Format format) {
    cudaEvent_t start, stop;
    cudaError_t err;
    void *devPtr, *hostPtr;
    Profiling profile = {0};
    profile.err = true;

    cudaEventCreate(&start);
    cudaEventCreate(&stop);
    if (mode == PINNED) {
        err = cudaMallocHost(&hostPtr, size);
        if (err != cudaSuccess) {
            fprintf(stderr, "! host pinned memory allocation failed.\n! %s\n", 
                            cudaGetErrorString(err));
            return profile;
        }
    }else if (mode == PAGEABLE) {
        hostPtr = malloc(size);
        if (hostPtr == NULL) {
            printf("! host memory allocation failed!\n");
            return profile; // lets continue, though it is most likely not working
        }
    }else {
        return profile;
    }

    if (format == HUMAN) {
        printf(" > allocating %s%d%s bytes in device memory\n", YELLOW, size, WHITE);
    }
    cudaEventRecord(start);
    err = cudaMalloc(&devPtr, size);
    if (err != cudaSuccess) {
        fprintf(stderr, "! cuda error: %s\n", cudaGetErrorString(err));
        return profile;
    }

    cudaEventRecord(stop);
    cudaEventSynchronize(stop);
    cudaEventElapsedTime(&profile.mseconds_h2d, start, stop);
    if (format == HUMAN) {
        printf(" > allocated in %.2f milliseconds\n");
    }

    //Host --> Device
    cudaEventRecord(start);
    cudaMemcpy(devPtr, hostPtr, size, cudaMemcpyDefault);
    cudaEventRecord(stop);
    cudaEventSynchronize(stop);
    cudaEventElapsedTime(&profile.mseconds_h2d, start, stop);
    profile.bandwidth_h2d = calcBandwidth(size, profile.mseconds_h2d);
    if (format == HUMAN) {
        printf("  > copied host->device in %.2f milliseconds\n", profile.mseconds_h2d);
        printf("  > this equals a bandwidth of %s%.2f%s GB/s\n", RED, profile.bandwidth_h2d, WHITE);
    }
    
    // Device --> Host
    cudaEventRecord(start);
    cudaMemcpy(hostPtr, devPtr, size, cudaMemcpyDefault);
    cudaEventRecord(stop);
    cudaEventSynchronize(stop);
    cudaEventElapsedTime(&profile.mseconds_d2h, start, stop);
    profile.bandwidth_d2h = calcBandwidth(size, profile.mseconds_d2h);
    if (format == HUMAN) {
        printf("  > copied device->host in %.2f milliseconds\n", profile.mseconds_d2h);
        printf("  > this equals a bandwidth of %.2f GB/s\n", profile.bandwidth_d2h);
    }

    if (mode == PINNED) {
        cudaFreeHost(hostPtr);
    }else if (mode == PAGEABLE) {
        free(hostPtr);
    }
    cudaFree(devPtr);

    profile.err = false;
    return profile;
}


void showUsage(void) {
    printf("""usage: bandwidth [-p] [-o]\n""");
}



int main(int argc, char** argv)
{
    int num_samples = 0; // number of samples
    Mode mode = PAGEABLE; // Transfer-mode, pinned or pageable
    Format format = HUMAN; // Format: human-readable or plotable
    bool verbose = false;
    char *output = NULL;
    FILE *fd = NULL;

    while(1) {
        int option_index = 0;
        int c = getopt_long(argc, argv, "n:m:f:ho:", long_options, &option_index);
        if (c == -1) {
            break;
        }
        switch(c) {
            case 'm':
                if (optarg) {
                    if (strcmp(optarg, "pinned") == 0) {
                        mode = PINNED;
                    } else if ( strcmp(optarg, "pageable") == 0) {
                        mode = PAGEABLE;
                    }
                }
            break;

            case 'f':
                if (optarg) {
                    if (strcmp(optarg, "plot") == 0) {
                        format = PLOT;
                    } else if (strcmp(optarg, "human") == 0 ||
                               strcmp(optarg, "plain") == 0) {
                        format = HUMAN;
                    } 
                }
                break;

            case 'o': //output file
                if (optarg) {
                    size_t len = strnlen(optarg, 255);
                    output = (char*) malloc(sizeof(char) * len + 1);
                    strncpy(output, optarg, len);
                    output[len] = 0; // null-terminate
                }
                break;
            case 'h':
                showUsage();
                exit(0);

            case 'n': // --num
                if (optarg) {
                    num_samples = atoi(optarg);
                }
                break;

            case 'v': // verbose
                verbose = true;
                break;


            default:
                printf("error\n");
                exit(-1);
        }
    }
    if (num_samples != 0 && format == HUMAN) {
        fprintf(stderr, "!! can't provide amount of samples for human format.\n");
        exit(1);
    }
    cudaDeviceReset();

    size_t *samples = NULL;
    if (format == HUMAN) {
        num_samples = 0;
        size_t samples_tmp[] = {1, 64, 512, 1024, 1024*1024, 1024*1024*10, 1024*1024*10, 1024*1024*200, 1024*1024*500, 1024*1024*900};
        num_samples = sizeof(samples_tmp) / sizeof(samples_tmp[0]);
        printf("num: %d\n", num_samples);
        samples =(size_t*) malloc(sizeof(size_t) * num_samples);
        for(int i=0; i < num_samples; i ++) {
            *(samples+i) = samples_tmp[i]; // i know, i know...
        } // could've used C++ new op...
    }else if (format == PLOT) { // use more data points for plotting
        num_samples = num_samples ? num_samples :  128;
        samples = (size_t*)malloc(sizeof(size_t) * num_samples);
        size_t max_size = 1024*1024*700;
        int i=0;
        for (size_t size = 1; size < max_size; size += max_size/num_samples, i++) {
            samples[i] = size;
        }
    }


    Profiling profile = {0};
    if (output != NULL) {
        fd = fopen(output, "a+");
        if (fd == NULL) {
            fprintf(stderr, "!! error opening output-file.\n");
            exit(1);
        }
    }else {
        fd = stdout;
    }
    for (int i=0; i < num_samples; i++) {
        size_t size = samples[i];
        if (format == PLOT) {
            if (i % (num_samples/10) == 0 && i != 0) {
                fprintf(stderr, "%d[%d]\n", i, size);
            }else {
                fprintf(stderr, ".");
            }
        }

        if (format == HUMAN) {
            printf("* test #%d: %d bytes ( %.2f MB)", i, size, size / (float)(1024*1024));
            printf("> allocating %d bytes in host memory\n", size);
        }
        profile = profileTransfer(size, mode, format);
        fprintf(fd, "%d\t%f\t%f\t%f\n", size, profile.bandwidth_h2d, profile.bandwidth_d2h,
                                        (profile.bandwidth_h2d + profile.bandwidth_d2h) / 2.0);
    }
    if (fd != NULL) {
        fclose(fd);
    }
    free(samples);

    return 0;
}


	.file	"tmpxft_0027e204_00000000-6_bandwidth.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2033:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2033:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z13calcBandwidthmf
	.type	_Z13calcBandwidthmf, @function
_Z13calcBandwidthmf:
.LFB2027:
	.cfi_startproc
	endbr64
	testq	%rdi, %rdi
	js	.L3
	cvtsi2ssq	%rdi, %xmm1
	jmp	.L4
.L3:
	movq	%rdi, %rax
	andl	$1, %edi
	shrq	%rax
	orq	%rdi, %rax
	cvtsi2ssq	%rax, %xmm1
	addss	%xmm1, %xmm1
.L4:
	mulss	.LC1(%rip), %xmm1
	divss	.LC2(%rip), %xmm0
	divss	%xmm0, %xmm1
	movaps	%xmm1, %xmm0
	ret
	.cfi_endproc
.LFE2027:
	.size	_Z13calcBandwidthmf, .-_Z13calcBandwidthmf
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC3:
	.string	"! host pinned memory allocation failed.\n! %s\n"
.LC4:
	.string	"! host memory allocation failed!\n"
.LC5:
	.string	"\033[1;00m"
.LC6:
	.string	"\033[1;33m"
.LC7:
	.string	" > allocating %s%d%s bytes in device memory\n"
.LC8:
	.string	"! cuda error: %s\n"
.LC9:
	.string	" > allocated in %.2f milliseconds\n"
.LC10:
	.string	"  > copied host->device in %.2f milliseconds\n"
.LC11:
	.string	"\033[1;31m"
.LC12:
	.string	"  > this equals a bandwidth of %s%.2f%s GB/s\n"
.LC13:
	.string	"  > copied device->host in %.2f milliseconds\n"
.LC14:
	.string	"  > this equals a bandwidth of %.2f GB/s\n"
	.text
	.globl	_Z15profileTransferm4Mode6Format
	.type	_Z15profileTransferm4Mode6Format, @function
_Z15profileTransferm4Mode6Format:
.LFB2028:
	.cfi_startproc
	endbr64
	pushq	%r14
	.cfi_def_cfa_offset 16
	.cfi_offset 14, -16
	pushq	%r13
	.cfi_def_cfa_offset 24
	.cfi_offset 13, -24
	movl	%edx, %r13d
	pushq	%r12
	.cfi_def_cfa_offset 32
	.cfi_offset 12, -32
	movl	%ecx, %r12d
	movl	$5, %ecx
	pushq	%rbp
	.cfi_def_cfa_offset 40
	.cfi_offset 6, -40
	movq	%rsi, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset 3, -48
	movq	%rdi, %rbx
	subq	$48, %rsp
	.cfi_def_cfa_offset 96
	movq	%fs:40, %rax
	movq	%rax, 40(%rsp)
	xorl	%eax, %eax
	rep stosl
	leaq	8(%rsp), %rdi
	movb	$1, 16(%rbx)
	call	cudaEventCreate@PLT
	leaq	16(%rsp), %rdi
	call	cudaEventCreate@PLT
	testl	%r13d, %r13d
	jne	.L7
	leaq	32(%rsp), %rdi
	movq	%rbp, %rsi
	call	cudaMallocHost@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L8
	call	cudaGetErrorString@PLT
	leaq	.LC3(%rip), %rdx
	movq	%rax, %rcx
	jmp	.L24
.L7:
	cmpl	$1, %r13d
	jne	.L6
	movq	%rbp, %rdi
	call	malloc@PLT
	movq	%rax, 32(%rsp)
	testq	%rax, %rax
	jne	.L8
	leaq	.LC4(%rip), %rsi
	movl	$2, %edi
	call	__printf_chk@PLT
	jmp	.L6
.L8:
	cmpl	$1, %r12d
	jne	.L10
	leaq	.LC5(%rip), %r8
	movq	%rbp, %rcx
	xorl	%eax, %eax
	movl	$2, %edi
	leaq	.LC6(%rip), %rdx
	leaq	.LC7(%rip), %rsi
	call	__printf_chk@PLT
.L10:
	movq	8(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	leaq	24(%rsp), %rdi
	movq	%rbp, %rsi
	call	cudaMalloc@PLT
	movl	%eax, %edi
	testl	%eax, %eax
	je	.L11
	call	cudaGetErrorString@PLT
	leaq	.LC8(%rip), %rdx
	movq	%rax, %rcx
.L24:
	movq	stderr(%rip), %rdi
	movl	$2, %esi
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
	jmp	.L6
.L11:
	movq	16(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movq	16(%rsp), %rdi
	call	cudaEventSynchronize@PLT
	movq	16(%rsp), %rdx
	movq	8(%rsp), %rsi
	movq	%rbx, %rdi
	call	cudaEventElapsedTime@PLT
	cmpl	$1, %r12d
	jne	.L12
	leaq	.LC9(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
.L12:
	movq	8(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movq	32(%rsp), %rsi
	movq	24(%rsp), %rdi
	movq	%rbp, %rdx
	movl	$4, %ecx
	call	cudaMemcpy@PLT
	movq	16(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movq	16(%rsp), %rdi
	call	cudaEventSynchronize@PLT
	movq	16(%rsp), %rdx
	movq	8(%rsp), %rsi
	movq	%rbx, %rdi
	call	cudaEventElapsedTime@PLT
	movss	(%rbx), %xmm0
	testq	%rbp, %rbp
	js	.L13
	cvtsi2ssq	%rbp, %xmm1
	jmp	.L14
.L13:
	movq	%rbp, %rax
	movq	%rbp, %rdx
	shrq	%rax
	andl	$1, %edx
	orq	%rdx, %rax
	cvtsi2ssq	%rax, %xmm1
	addss	%xmm1, %xmm1
.L14:
	mulss	.LC1(%rip), %xmm1
	movaps	%xmm0, %xmm2
	divss	.LC2(%rip), %xmm2
	movd	%xmm1, %r14d
	divss	%xmm2, %xmm1
	movss	%xmm1, 8(%rbx)
	cmpl	$1, %r12d
	jne	.L15
	leaq	.LC10(%rip), %rsi
	movl	$2, %edi
	movb	$1, %al
	cvtss2sd	%xmm0, %xmm0
	call	__printf_chk@PLT
	leaq	.LC5(%rip), %rcx
	movl	$2, %edi
	leaq	.LC11(%rip), %rdx
	leaq	.LC12(%rip), %rsi
	movb	$1, %al
	cvtss2sd	8(%rbx), %xmm0
	call	__printf_chk@PLT
.L15:
	movq	8(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movq	24(%rsp), %rsi
	movq	32(%rsp), %rdi
	movq	%rbp, %rdx
	movl	$4, %ecx
	call	cudaMemcpy@PLT
	movq	16(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movq	16(%rsp), %rdi
	call	cudaEventSynchronize@PLT
	movq	16(%rsp), %rdx
	movq	8(%rsp), %rsi
	leaq	4(%rbx), %rdi
	call	cudaEventElapsedTime@PLT
	movss	4(%rbx), %xmm0
	movd	%r14d, %xmm1
	decl	%r12d
	movaps	%xmm0, %xmm2
	divss	.LC2(%rip), %xmm2
	divss	%xmm2, %xmm1
	movss	%xmm1, 12(%rbx)
	jne	.L16
	leaq	.LC13(%rip), %rsi
	movl	$2, %edi
	movb	$1, %al
	cvtss2sd	%xmm0, %xmm0
	call	__printf_chk@PLT
	leaq	.LC14(%rip), %rsi
	movl	$2, %edi
	movb	$1, %al
	cvtss2sd	12(%rbx), %xmm0
	call	__printf_chk@PLT
.L16:
	movq	32(%rsp), %rdi
	testl	%r13d, %r13d
	jne	.L17
	call	cudaFreeHost@PLT
	jmp	.L18
.L17:
	call	free@PLT
.L18:
	movq	24(%rsp), %rdi
	call	cudaFree@PLT
	movb	$0, 16(%rbx)
.L6:
	movq	40(%rsp), %rax
	subq	%fs:40, %rax
	je	.L19
	call	__stack_chk_fail@PLT
.L19:
	addq	$48, %rsp
	.cfi_def_cfa_offset 48
	movq	%rbx, %rax
	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
.LFE2028:
	.size	_Z15profileTransferm4Mode6Format, .-_Z15profileTransferm4Mode6Format
	.section	.rodata.str1.1
.LC15:
	.string	"usage: bandwidth [-p] [-o]\n"
	.text
	.globl	_Z9showUsagev
	.type	_Z9showUsagev, @function
_Z9showUsagev:
.LFB2029:
	.cfi_startproc
	endbr64
	leaq	.LC15(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	jmp	__printf_chk@PLT
	.cfi_endproc
.LFE2029:
	.size	_Z9showUsagev, .-_Z9showUsagev
	.section	.rodata.str1.1
.LC16:
	.string	"n:m:f:ho:"
.LC17:
	.string	"pinned"
.LC18:
	.string	"pageable"
.LC19:
	.string	"plot"
.LC20:
	.string	"human"
.LC21:
	.string	"plain"
.LC22:
	.string	"error\n"
.LC23:
	.string	"!! can't provide amount of samples for human format.\n"
.LC24:
	.string	"num: %d\n"
.LC25:
	.string	"a+"
.LC26:
	.string	"!! error opening output-file.\n"
.LC28:
	.string	"* test #%d: %d bytes ( %.2f MB)"
.LC29:
	.string	"> allocating %d bytes in host memory\n"
.LC30:
	.string	"%d[%d]\n"
.LC31:
	.string	"."
.LC33:
	.string	"%d\t%f\t%f\t%f\n"
	.section	.rodata
	.align 32
.LC0:
	.quad	1
	.quad	64
	.quad	512
	.quad	1024
	.quad	1048576
	.quad	10485760
	.quad	10485760
	.quad	209715200
	.quad	524288000
	.quad	943718400
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2030:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movl	$1, %r15d
	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
	movl	$1, %r12d
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	xorl	%ebp, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	xorl	%ebx, %ebx
	subq	$152, %rsp
	.cfi_def_cfa_offset 208
	movl	%edi, 8(%rsp)
	movq	%rsi, 24(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 136(%rsp)
	xorl	%eax, %eax
.L36:
	xorl	%eax, %eax
	movq	24(%rsp), %rsi
	movl	8(%rsp), %edi
	leaq	long_options(%rip), %rcx
	movl	%eax, 36(%rsp)
	leaq	36(%rsp), %rax
	leaq	.LC16(%rip), %rdx
	movq	%rax, %r8
	movq	%rax, 16(%rsp)
	call	getopt_long@PLT
	cmpl	$-1, %eax
	je	.L27
	subl	$102, %eax
	cmpl	$16, %eax
	ja	.L28
	leaq	.L30(%rip), %rdx
	movslq	(%rdx,%rax,4), %rax
	addq	%rdx, %rax
	notrack jmp	*%rax
	.section	.rodata
	.align 4
	.align 4
.L30:
	.long	.L35-.L30
	.long	.L28-.L30
	.long	.L34-.L30
	.long	.L28-.L30
	.long	.L28-.L30
	.long	.L28-.L30
	.long	.L28-.L30
	.long	.L33-.L30
	.long	.L32-.L30
	.long	.L31-.L30
	.long	.L28-.L30
	.long	.L28-.L30
	.long	.L28-.L30
	.long	.L28-.L30
	.long	.L28-.L30
	.long	.L28-.L30
	.long	.L36-.L30
	.section	.text.startup
.L33:
	movq	optarg(%rip), %r13
	testq	%r13, %r13
	je	.L36
	leaq	.LC17(%rip), %rsi
	movq	%r13, %rdi
	call	strcmp@PLT
	testl	%eax, %eax
	je	.L53
	leaq	.LC18(%rip), %rsi
	movq	%r13, %rdi
	call	strcmp@PLT
	testl	%eax, %eax
	movl	$1, %eax
	cmove	%eax, %r15d
	jmp	.L36
.L35:
	movq	optarg(%rip), %r13
	testq	%r13, %r13
	je	.L36
	leaq	.LC19(%rip), %rsi
	movq	%r13, %rdi
	call	strcmp@PLT
	testl	%eax, %eax
	je	.L55
	leaq	.LC20(%rip), %rsi
	movq	%r13, %rdi
	call	strcmp@PLT
	testl	%eax, %eax
	je	.L56
	leaq	.LC21(%rip), %rsi
	movq	%r13, %rdi
	call	strcmp@PLT
	testl	%eax, %eax
	movl	$1, %eax
	cmove	%eax, %r12d
	jmp	.L36
.L31:
	movq	optarg(%rip), %r14
	testq	%r14, %r14
	je	.L36
	movl	$255, %esi
	movq	%r14, %rdi
	call	strnlen@PLT
	leaq	1(%rax), %rcx
	movq	%rax, %r13
	movq	%rcx, %rdi
	movq	%rcx, 16(%rsp)
	call	malloc@PLT
	movq	16(%rsp), %rcx
	movq	%r13, %rdx
	movq	%r14, %rsi
	movq	%rax, %rbp
	movq	%rax, %rdi
	call	__strncpy_chk@PLT
	movb	$0, 0(%rbp,%r13)
	jmp	.L36
.L34:
	call	_Z9showUsagev
	xorl	%edi, %edi
	jmp	.L87
.L32:
	movq	optarg(%rip), %rdi
	testq	%rdi, %rdi
	je	.L36
	call	atoi@PLT
	movl	%eax, %ebx
	jmp	.L36
.L28:
	movl	$2, %edi
	leaq	.LC22(%rip), %rsi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	orl	$-1, %edi
.L87:
	call	exit@PLT
.L53:
	xorl	%r15d, %r15d
	jmp	.L36
.L55:
	xorl	%r12d, %r12d
	jmp	.L36
.L56:
	movl	$1, %r12d
	jmp	.L36
.L27:
	testl	%ebx, %ebx
	je	.L37
	leaq	.LC23(%rip), %rdx
	testl	$1, %r12d
	jne	.L88
.L37:
	call	cudaDeviceReset@PLT
	cmpl	$1, %r12d
	jne	.L38
	movl	$10, %edx
	leaq	56(%rsp), %rdi
	leaq	.LC0(%rip), %rsi
	xorl	%eax, %eax
	movl	$20, %ecx
	movl	$10, %ebx
	rep movsl
	leaq	.LC24(%rip), %rsi
	movl	$2, %edi
	call	__printf_chk@PLT
	movl	$80, %edi
	call	malloc@PLT
	leaq	56(%rsp), %rsi
	movl	$20, %ecx
	movq	%rax, %rdi
	movq	%rax, 8(%rsp)
	rep movsl
	jmp	.L39
.L38:
	testl	%ebx, %ebx
	jne	.L40
	movl	$128, %ebx
.L40:
	movslq	%ebx, %r13
	leaq	0(,%r13,8), %rdi
	call	malloc@PLT
	xorl	%edx, %edx
	movq	%rax, 8(%rsp)
	movl	$734003200, %eax
	movq	8(%rsp), %rcx
	divq	%r13
	movl	$1, %edx
.L41:
	movq	%rdx, (%rcx)
	addq	%rax, %rdx
	addq	$8, %rcx
	cmpq	$734003199, %rdx
	jbe	.L41
.L39:
	testq	%rbp, %rbp
	je	.L42
	leaq	.LC25(%rip), %rsi
	movq	%rbp, %rdi
	call	fopen@PLT
	movq	%rax, %r13
	testq	%rax, %rax
	jne	.L43
	leaq	.LC26(%rip), %rdx
.L88:
	movq	stderr(%rip), %rdi
	movl	$2, %esi
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
	movl	$1, %edi
	jmp	.L87
.L42:
	movq	stdout(%rip), %r13
.L43:
	movl	%ebx, %eax
	movl	$10, %ecx
	xorl	%r14d, %r14d
	cltd
	idivl	%ecx
	movl	%eax, 24(%rsp)
.L44:
	movl	%r14d, %esi
	cmpl	%r14d, %ebx
	jle	.L89
	movq	8(%rsp), %rax
	movq	(%rax,%r14,8), %rbp
	testl	%r12d, %r12d
	je	.L45
	testq	%rbp, %rbp
	js	.L46
	cvtsi2ssq	%rbp, %xmm0
	jmp	.L47
.L46:
	movq	%rbp, %rax
	movq	%rbp, %rdx
	shrq	%rax
	andl	$1, %edx
	orq	%rdx, %rax
	cvtsi2ssq	%rax, %xmm0
	addss	%xmm0, %xmm0
.L47:
	mulss	.LC27(%rip), %xmm0
	movl	%esi, %edx
	movq	%rbp, %rcx
	movb	$1, %al
	leaq	.LC28(%rip), %rsi
	movl	$2, %edi
	cvtss2sd	%xmm0, %xmm0
	call	__printf_chk@PLT
	movq	%rbp, %rdx
	movl	$2, %edi
	leaq	.LC29(%rip), %rsi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	jmp	.L48
.L45:
	movl	%esi, %eax
	movq	stderr(%rip), %rdi
	cltd
	idivl	24(%rsp)
	testl	%edx, %edx
	jne	.L49
	testl	%esi, %esi
	je	.L49
	movl	%esi, %ecx
	movq	%rbp, %r8
	leaq	.LC30(%rip), %rdx
	xorl	%eax, %eax
	movl	$2, %esi
	call	__fprintf_chk@PLT
	jmp	.L48
.L49:
	leaq	.LC31(%rip), %rdx
	movl	$2, %esi
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
.L48:
	movq	16(%rsp), %rdi
	movl	%r12d, %ecx
	movl	%r15d, %edx
	movq	%rbp, %rsi
	incq	%r14
	call	_Z15profileTransferm4Mode6Format
	movss	44(%rsp), %xmm0
	movq	%rbp, %rcx
	movb	$3, %al
	movss	48(%rsp), %xmm1
	leaq	.LC33(%rip), %rdx
	movl	$2, %esi
	movq	%r13, %rdi
	movaps	%xmm0, %xmm2
	cvtss2sd	%xmm0, %xmm0
	addss	%xmm1, %xmm2
	cvtss2sd	%xmm1, %xmm1
	cvtss2sd	%xmm2, %xmm2
	mulsd	.LC32(%rip), %xmm2
	call	__fprintf_chk@PLT
	jmp	.L44
.L89:
	testq	%r13, %r13
	je	.L51
	movq	%r13, %rdi
	call	fclose@PLT
.L51:
	movq	8(%rsp), %rdi
	call	free@PLT
	movq	136(%rsp), %rax
	subq	%fs:40, %rax
	je	.L52
	call	__stack_chk_fail@PLT
.L52:
	addq	$152, %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
.LFE2030:
	.size	main, .-main
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2056:
	.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
.LFE2056:
	.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
	.globl	long_options
	.section	.rodata.str1.1
.LC35:
	.string	"num"
.LC36:
	.string	"mode"
.LC37:
	.string	"format"
.LC38:
	.string	"help"
	.section	.data.rel.local,"aw"
	.align 32
	.type	long_options, @object
	.size	long_options, 160
long_options:
	.quad	.LC35
	.long	1
	.zero	4
	.quad	0
	.long	0
	.zero	4
	.quad	.LC36
	.long	1
	.zero	4
	.quad	0
	.long	0
	.zero	4
	.quad	.LC37
	.long	1
	.zero	4
	.quad	0
	.long	0
	.zero	4
	.quad	.LC38
	.long	0
	.zero	4
	.quad	0
	.long	0
	.zero	4
	.quad	0
	.long	0
	.zero	4
	.quad	0
	.long	0
	.zero	4
	.section	.rodata.cst4,"aM",@progbits,4
	.align 4
.LC1:
	.long	813694976
	.align 4
.LC2:
	.long	1148846080
	.align 4
.LC27:
	.long	897581056
	.section	.rodata.cst8,"aM",@progbits,8
	.align 8
.LC32:
	.long	0
	.long	1071644672
	.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:

/*
 * bandwidth-test, used for plotting
 * can use pageable or pinned memory
 */
#include <stdio.h>
#include <getopt.h>
#include <hip/hip_runtime.h>

#define YELLOW "\e[1;33m"
#define RED    "\e[1;31m"
#define WHITE  "\e[1;00m"



typedef struct {
    float mseconds_h2d;
    float mseconds_d2h;
    float bandwidth_h2d;
    float bandwidth_d2h;
    bool err;
} Profiling;

typedef enum {
    PLOT,
    HUMAN
} Format;

typedef enum {
    PINNED,
    PAGEABLE
} Mode;



struct option long_options[] = {
    {"num"      , required_argument, 0, 0},
    {"mode"     , required_argument, 0, 0},
    {"format"   , required_argument, 0, 0},
    {"help"     , 0                , 0, 0},
    {0          , 0                , 0, 0}
};

float calcBandwidth(size_t size, float mseconds) {
    return (size/((float)1024*1024*1024)) / (mseconds / 1000.0f);
}

/*! perform a profiled memory transfer
 *
 *  with pinned or pageable memory using size bytes
 *
 */
Profiling profileTransfer(size_t size, Mode mode, Format format) {
    hipEvent_t start, stop;
    hipError_t err;
    void *devPtr, *hostPtr;
    Profiling profile = {0};
    profile.err = true;

    hipEventCreate(&start);
    hipEventCreate(&stop);
    if (mode == PINNED) {
        err = hipHostMalloc(&hostPtr, size, hipHostMallocDefault);
        if (err != hipSuccess) {
            fprintf(stderr, "! host pinned memory allocation failed.\n! %s\n", 
                            hipGetErrorString(err));
            return profile;
        }
    }else if (mode == PAGEABLE) {
        hostPtr = malloc(size);
        if (hostPtr == NULL) {
            printf("! host memory allocation failed!\n");
            return profile; // lets continue, though it is most likely not working
        }
    }else {
        return profile;
    }

    if (format == HUMAN) {
        printf(" > allocating %s%d%s bytes in device memory\n", YELLOW, size, WHITE);
    }
    hipEventRecord(start);
    err = hipMalloc(&devPtr, size);
    if (err != hipSuccess) {
        fprintf(stderr, "! cuda error: %s\n", hipGetErrorString(err));
        return profile;
    }

    hipEventRecord(stop);
    hipEventSynchronize(stop);
    hipEventElapsedTime(&profile.mseconds_h2d, start, stop);
    if (format == HUMAN) {
        printf(" > allocated in %.2f milliseconds\n");
    }

    //Host --> Device
    hipEventRecord(start);
    hipMemcpy(devPtr, hostPtr, size, hipMemcpyDefault);
    hipEventRecord(stop);
    hipEventSynchronize(stop);
    hipEventElapsedTime(&profile.mseconds_h2d, start, stop);
    profile.bandwidth_h2d = calcBandwidth(size, profile.mseconds_h2d);
    if (format == HUMAN) {
        printf("  > copied host->device in %.2f milliseconds\n", profile.mseconds_h2d);
        printf("  > this equals a bandwidth of %s%.2f%s GB/s\n", RED, profile.bandwidth_h2d, WHITE);
    }
    
    // Device --> Host
    hipEventRecord(start);
    hipMemcpy(hostPtr, devPtr, size, hipMemcpyDefault);
    hipEventRecord(stop);
    hipEventSynchronize(stop);
    hipEventElapsedTime(&profile.mseconds_d2h, start, stop);
    profile.bandwidth_d2h = calcBandwidth(size, profile.mseconds_d2h);
    if (format == HUMAN) {
        printf("  > copied device->host in %.2f milliseconds\n", profile.mseconds_d2h);
        printf("  > this equals a bandwidth of %.2f GB/s\n", profile.bandwidth_d2h);
    }

    if (mode == PINNED) {
        hipHostFree(hostPtr);
    }else if (mode == PAGEABLE) {
        free(hostPtr);
    }
    hipFree(devPtr);

    profile.err = false;
    return profile;
}


void showUsage(void) {
    printf("""usage: bandwidth [-p] [-o]\n""");
}



int main(int argc, char** argv)
{
    int num_samples = 0; // number of samples
    Mode mode = PAGEABLE; // Transfer-mode, pinned or pageable
    Format format = HUMAN; // Format: human-readable or plotable
    bool verbose = false;
    char *output = NULL;
    FILE *fd = NULL;

    while(1) {
        int option_index = 0;
        int c = getopt_long(argc, argv, "n:m:f:ho:", long_options, &option_index);
        if (c == -1) {
            break;
        }
        switch(c) {
            case 'm':
                if (optarg) {
                    if (strcmp(optarg, "pinned") == 0) {
                        mode = PINNED;
                    } else if ( strcmp(optarg, "pageable") == 0) {
                        mode = PAGEABLE;
                    }
                }
            break;

            case 'f':
                if (optarg) {
                    if (strcmp(optarg, "plot") == 0) {
                        format = PLOT;
                    } else if (strcmp(optarg, "human") == 0 ||
                               strcmp(optarg, "plain") == 0) {
                        format = HUMAN;
                    } 
                }
                break;

            case 'o': //output file
                if (optarg) {
                    size_t len = strnlen(optarg, 255);
                    output = (char*) malloc(sizeof(char) * len + 1);
                    strncpy(output, optarg, len);
                    output[len] = 0; // null-terminate
                }
                break;
            case 'h':
                showUsage();
                exit(0);

            case 'n': // --num
                if (optarg) {
                    num_samples = atoi(optarg);
                }
                break;

            case 'v': // verbose
                verbose = true;
                break;


            default:
                printf("error\n");
                exit(-1);
        }
    }
    if (num_samples != 0 && format == HUMAN) {
        fprintf(stderr, "!! can't provide amount of samples for human format.\n");
        exit(1);
    }
    hipDeviceReset();

    size_t *samples = NULL;
    if (format == HUMAN) {
        num_samples = 0;
        size_t samples_tmp[] = {1, 64, 512, 1024, 1024*1024, 1024*1024*10, 1024*1024*10, 1024*1024*200, 1024*1024*500, 1024*1024*900};
        num_samples = sizeof(samples_tmp) / sizeof(samples_tmp[0]);
        printf("num: %d\n", num_samples);
        samples =(size_t*) malloc(sizeof(size_t) * num_samples);
        for(int i=0; i < num_samples; i ++) {
            *(samples+i) = samples_tmp[i]; // i know, i know...
        } // could've used C++ new op...
    }else if (format == PLOT) { // use more data points for plotting
        num_samples = num_samples ? num_samples :  128;
        samples = (size_t*)malloc(sizeof(size_t) * num_samples);
        size_t max_size = 1024*1024*700;
        int i=0;
        for (size_t size = 1; size < max_size; size += max_size/num_samples, i++) {
            samples[i] = size;
        }
    }


    Profiling profile = {0};
    if (output != NULL) {
        fd = fopen(output, "a+");
        if (fd == NULL) {
            fprintf(stderr, "!! error opening output-file.\n");
            exit(1);
        }
    }else {
        fd = stdout;
    }
    for (int i=0; i < num_samples; i++) {
        size_t size = samples[i];
        if (format == PLOT) {
            if (i % (num_samples/10) == 0 && i != 0) {
                fprintf(stderr, "%d[%d]\n", i, size);
            }else {
                fprintf(stderr, ".");
            }
        }

        if (format == HUMAN) {
            printf("* test #%d: %d bytes ( %.2f MB)", i, size, size / (float)(1024*1024));
            printf("> allocating %d bytes in host memory\n", size);
        }
        profile = profileTransfer(size, mode, format);
        fprintf(fd, "%d\t%f\t%f\t%f\n", size, profile.bandwidth_h2d, profile.bandwidth_d2h,
                                        (profile.bandwidth_h2d + profile.bandwidth_d2h) / 2.0);
    }
    if (fd != NULL) {
        fclose(fd);
    }
    free(samples);

    return 0;
}


	.text
	.file	"bandwidth.hip"
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0                          # -- Begin function _Z13calcBandwidthmf
.LCPI0_0:
	.long	0x30800000                      # float 9.31322574E-10
.LCPI0_1:
	.long	0x447a0000                      # float 1000
	.text
	.globl	_Z13calcBandwidthmf
	.type	_Z13calcBandwidthmf,@function
_Z13calcBandwidthmf:                    # @_Z13calcBandwidthmf
	.cfi_startproc
# %bb.0:
	testq	%rdi, %rdi
	js	.LBB0_1
# %bb.2:
	cvtsi2ss	%rdi, %xmm1
	jmp	.LBB0_3
.LBB0_1:
	movq	%rdi, %rax
	shrq	%rax
	andl	$1, %edi
	orq	%rax, %rdi
	cvtsi2ss	%rdi, %xmm1
	addss	%xmm1, %xmm1
.LBB0_3:
	mulss	.LCPI0_0(%rip), %xmm1
	divss	.LCPI0_1(%rip), %xmm0
	divss	%xmm0, %xmm1
	movaps	%xmm1, %xmm0
	retq
.Lfunc_end0:
	.size	_Z13calcBandwidthmf, .Lfunc_end0-_Z13calcBandwidthmf
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0                          # -- Begin function _Z15profileTransferm4Mode6Format
.LCPI1_0:
	.long	0x30800000                      # float 9.31322574E-10
.LCPI1_1:
	.long	0x447a0000                      # float 1000
	.text
	.globl	_Z15profileTransferm4Mode6Format
	.type	_Z15profileTransferm4Mode6Format,@function
_Z15profileTransferm4Mode6Format:       # @_Z15profileTransferm4Mode6Format
	.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	%rbx
	.cfi_def_cfa_offset 40
	subq	$40, %rsp
	.cfi_def_cfa_offset 80
	.cfi_offset %rbx, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movl	%ecx, %r14d
	movl	%edx, %ebp
	movq	%rsi, %r15
	movq	%rdi, %rbx
	xorps	%xmm0, %xmm0
	movups	%xmm0, (%rdi)
	movl	$0, 16(%rdi)
	movb	$1, 16(%rdi)
	leaq	8(%rsp), %rdi
	callq	hipEventCreate
	movq	%rsp, %rdi
	callq	hipEventCreate
	testl	%ebp, %ebp
	je	.LBB1_1
# %bb.4:
	cmpl	$1, %ebp
	jne	.LBB1_24
# %bb.5:
	movq	%r15, %rdi
	callq	malloc
	movq	%rax, 16(%rsp)
	testq	%rax, %rax
	jne	.LBB1_7
# %bb.6:
	movl	$.Lstr, %edi
	callq	puts@PLT
	jmp	.LBB1_24
.LBB1_1:
	leaq	16(%rsp), %rdi
	movq	%r15, %rsi
	xorl	%edx, %edx
	callq	hipHostMalloc
	testl	%eax, %eax
	jne	.LBB1_2
.LBB1_7:
	cmpl	$1, %r14d
	jne	.LBB1_9
# %bb.8:
	movl	$.L.str.6, %edi
	movl	$.L.str.7, %esi
	movl	$.L.str.8, %ecx
	movq	%r15, %rdx
	xorl	%eax, %eax
	callq	printf
.LBB1_9:
	movq	8(%rsp), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	leaq	24(%rsp), %rdi
	movq	%r15, %rsi
	callq	hipMalloc
	testl	%eax, %eax
	jne	.LBB1_10
# %bb.11:
	movq	(%rsp), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	(%rsp), %rdi
	callq	hipEventSynchronize
	movq	8(%rsp), %rsi
	movq	(%rsp), %rdx
	movq	%rbx, %rdi
	callq	hipEventElapsedTime
	cmpl	$1, %r14d
	jne	.LBB1_13
# %bb.12:
	movl	$.L.str.10, %edi
	xorl	%eax, %eax
	callq	printf
.LBB1_13:
	movq	8(%rsp), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	24(%rsp), %rdi
	movq	16(%rsp), %rsi
	movq	%r15, %rdx
	movl	$4, %ecx
	callq	hipMemcpy
	movq	(%rsp), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	(%rsp), %rdi
	callq	hipEventSynchronize
	movq	8(%rsp), %rsi
	movq	(%rsp), %rdx
	movq	%rbx, %rdi
	callq	hipEventElapsedTime
	testq	%r15, %r15
	js	.LBB1_14
# %bb.15:
	cvtsi2ss	%r15, %xmm2
	jmp	.LBB1_16
.LBB1_14:
	movq	%r15, %rax
	shrq	%rax
	movl	%r15d, %ecx
	andl	$1, %ecx
	orq	%rax, %rcx
	cvtsi2ss	%rcx, %xmm2
	addss	%xmm2, %xmm2
.LBB1_16:
	mulss	.LCPI1_0(%rip), %xmm2
	movss	(%rbx), %xmm0                   # xmm0 = mem[0],zero,zero,zero
	movaps	%xmm0, %xmm1
	divss	.LCPI1_1(%rip), %xmm1
	movss	%xmm2, 36(%rsp)                 # 4-byte Spill
	divss	%xmm1, %xmm2
	movss	%xmm2, 8(%rbx)
	cmpl	$1, %r14d
	jne	.LBB1_18
# %bb.17:
	cvtss2sd	%xmm0, %xmm0
	movl	$.L.str.11, %edi
	movb	$1, %al
	callq	printf
	xorps	%xmm0, %xmm0
	cvtss2sd	8(%rbx), %xmm0
	movl	$.L.str.12, %edi
	movl	$.L.str.13, %esi
	movl	$.L.str.8, %edx
	movb	$1, %al
	callq	printf
.LBB1_18:
	movq	8(%rsp), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	16(%rsp), %rdi
	movq	24(%rsp), %rsi
	movq	%r15, %rdx
	movl	$4, %ecx
	callq	hipMemcpy
	movq	(%rsp), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	movq	(%rsp), %rdi
	callq	hipEventSynchronize
	leaq	4(%rbx), %rdi
	movq	8(%rsp), %rsi
	movq	(%rsp), %rdx
	callq	hipEventElapsedTime
	movss	4(%rbx), %xmm0                  # xmm0 = mem[0],zero,zero,zero
	movaps	%xmm0, %xmm1
	divss	.LCPI1_1(%rip), %xmm1
	movss	36(%rsp), %xmm2                 # 4-byte Reload
                                        # xmm2 = mem[0],zero,zero,zero
	divss	%xmm1, %xmm2
	movss	%xmm2, 12(%rbx)
	cmpl	$1, %r14d
	jne	.LBB1_20
# %bb.19:
	cvtss2sd	%xmm0, %xmm0
	movl	$.L.str.14, %edi
	movb	$1, %al
	callq	printf
	xorps	%xmm0, %xmm0
	cvtss2sd	12(%rbx), %xmm0
	movl	$.L.str.15, %edi
	movb	$1, %al
	callq	printf
.LBB1_20:
	movq	16(%rsp), %rdi
	testl	%ebp, %ebp
	je	.LBB1_21
# %bb.22:
	callq	free
	jmp	.LBB1_23
.LBB1_21:
	callq	hipHostFree
.LBB1_23:
	movq	24(%rsp), %rdi
	callq	hipFree
	movb	$0, 16(%rbx)
.LBB1_24:
	movq	%rbx, %rax
	addq	$40, %rsp
	.cfi_def_cfa_offset 40
	popq	%rbx
	.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
.LBB1_10:
	.cfi_def_cfa_offset 80
	movq	stderr(%rip), %r14
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.9, %esi
	jmp	.LBB1_3
.LBB1_2:
	movq	stderr(%rip), %r14
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.4, %esi
.LBB1_3:
	movq	%r14, %rdi
	movq	%rax, %rdx
	xorl	%eax, %eax
	callq	fprintf
	jmp	.LBB1_24
.Lfunc_end1:
	.size	_Z15profileTransferm4Mode6Format, .Lfunc_end1-_Z15profileTransferm4Mode6Format
	.cfi_endproc
                                        # -- End function
	.globl	_Z9showUsagev                   # -- Begin function _Z9showUsagev
	.type	_Z9showUsagev,@function
_Z9showUsagev:                          # @_Z9showUsagev
	.cfi_startproc
# %bb.0:
	movl	$.Lstr.1, %edi
	jmp	puts@PLT                        # TAILCALL
.Lfunc_end2:
	.size	_Z9showUsagev, .Lfunc_end2-_Z9showUsagev
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0                          # -- Begin function main
.LCPI3_0:
	.long	0x35800000                      # float 9.53674316E-7
	.section	.rodata.cst8,"aM",@progbits,8
	.p2align	3, 0x0
.LCPI3_1:
	.quad	0x3fe0000000000000              # double 0.5
	.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	$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, %r12
	movl	%edi, %r13d
	xorl	%r14d, %r14d
	leaq	20(%rsp), %r15
	movl	$1, (%rsp)                      # 4-byte Folded Spill
	movl	$1, 4(%rsp)                     # 4-byte Folded Spill
	movl	$0, 8(%rsp)                     # 4-byte Folded Spill
.LBB3_1:                                # =>This Inner Loop Header: Depth=1
	movl	$0, 20(%rsp)
	movl	$.L.str.17, %edx
	movl	$long_options, %ecx
	movl	%r13d, %edi
	movq	%r12, %rsi
	movq	%r15, %r8
	callq	getopt_long
                                        # kill: def $eax killed $eax def $rax
	leal	-102(%rax), %ecx
	cmpl	$9, %ecx
	ja	.LBB3_2
# %bb.49:                               #   in Loop: Header=BB3_1 Depth=1
	jmpq	*.LJTI3_0(,%rcx,8)
.LBB3_8:                                #   in Loop: Header=BB3_1 Depth=1
	movq	optarg(%rip), %rbx
	testq	%rbx, %rbx
	je	.LBB3_1
# %bb.9:                                #   in Loop: Header=BB3_1 Depth=1
	movl	$.L.str.20, %esi
	movq	%rbx, %rdi
	callq	strcmp
	testl	%eax, %eax
	je	.LBB3_10
# %bb.11:                               #   in Loop: Header=BB3_1 Depth=1
	movl	$.L.str.21, %esi
	movq	%rbx, %rdi
	callq	strcmp
	testl	%eax, %eax
	je	.LBB3_13
# %bb.12:                               #   in Loop: Header=BB3_1 Depth=1
	movl	$.L.str.22, %esi
	movq	%rbx, %rdi
	callq	strcmp
	testl	%eax, %eax
	jne	.LBB3_1
.LBB3_13:                               #   in Loop: Header=BB3_1 Depth=1
	movl	$1, (%rsp)                      # 4-byte Folded Spill
	jmp	.LBB3_1
.LBB3_2:                                #   in Loop: Header=BB3_1 Depth=1
	cmpl	$118, %eax
	je	.LBB3_1
	jmp	.LBB3_3
.LBB3_17:                               #   in Loop: Header=BB3_1 Depth=1
	movq	optarg(%rip), %rdi
	testq	%rdi, %rdi
	je	.LBB3_1
# %bb.18:                               #   in Loop: Header=BB3_1 Depth=1
	callq	atoi
	movl	%eax, 8(%rsp)                   # 4-byte Spill
	jmp	.LBB3_1
.LBB3_5:                                #   in Loop: Header=BB3_1 Depth=1
	movq	optarg(%rip), %rbx
	testq	%rbx, %rbx
	je	.LBB3_1
# %bb.6:                                #   in Loop: Header=BB3_1 Depth=1
	movl	$.L.str.18, %esi
	movq	%rbx, %rdi
	callq	strcmp
	testl	%eax, %eax
	je	.LBB3_7
# %bb.19:                               #   in Loop: Header=BB3_1 Depth=1
	movl	$.L.str.19, %esi
	movq	%rbx, %rdi
	callq	strcmp
	testl	%eax, %eax
	movl	4(%rsp), %eax                   # 4-byte Reload
	movl	$1, %ecx
	cmovel	%ecx, %eax
	movl	%eax, 4(%rsp)                   # 4-byte Spill
	jmp	.LBB3_1
.LBB3_14:                               #   in Loop: Header=BB3_1 Depth=1
	movq	optarg(%rip), %rbx
	testq	%rbx, %rbx
	je	.LBB3_1
# %bb.15:                               #   in Loop: Header=BB3_1 Depth=1
	movl	$255, %esi
	movq	%rbx, %rdi
	callq	strnlen
	movq	%rax, %rbp
	leaq	1(%rax), %rdi
	callq	malloc
	movq	%rax, %r14
	movq	%rax, %rdi
	movq	%rbx, %rsi
	movq	%rbp, %rdx
	callq	strncpy
	movb	$0, (%r14,%rbp)
	jmp	.LBB3_1
.LBB3_7:                                #   in Loop: Header=BB3_1 Depth=1
	movl	$0, 4(%rsp)                     # 4-byte Folded Spill
	jmp	.LBB3_1
.LBB3_10:                               #   in Loop: Header=BB3_1 Depth=1
	movl	$0, (%rsp)                      # 4-byte Folded Spill
	jmp	.LBB3_1
.LBB3_3:
	cmpl	$-1, %eax
	jne	.LBB3_4
# %bb.20:
	movl	8(%rsp), %ebp                   # 4-byte Reload
	testl	%ebp, %ebp
	je	.LBB3_24
# %bb.21:
	cmpl	$1, (%rsp)                      # 4-byte Folded Reload
	je	.LBB3_22
.LBB3_24:
	callq	hipDeviceReset
	cmpl	$1, (%rsp)                      # 4-byte Folded Reload
	jne	.LBB3_26
# %bb.25:                               # %.loopexit.loopexit
	movl	$10, %ebx
	movl	$.L.str.25, %edi
	movl	$10, %esi
	xorl	%eax, %eax
	callq	printf
	movl	$80, %edi
	callq	malloc
	movq	%rax, %r15
	movl	$10, %ecx
	movl	$.L__const.main.samples_tmp, %esi
	movq	%rax, %rdi
	rep;movsq (%rsi), %es:(%rdi)
	jmp	.LBB3_28
.LBB3_26:
	testl	%ebp, %ebp
	movl	$128, %ebx
	cmovnel	%ebp, %ebx
	movslq	%ebx, %r12
	leaq	(,%r12,8), %rdi
	callq	malloc
	movq	%rax, %r15
	movl	$734003200, %eax                # imm = 0x2BC00000
	xorl	%edx, %edx
	divq	%r12
	movl	$1, %ecx
	movq	%r15, %rdx
.LBB3_27:                               # =>This Inner Loop Header: Depth=1
	movq	%rcx, (%rdx)
	addq	%rax, %rcx
	addq	$8, %rdx
	cmpq	$734003200, %rcx                # imm = 0x2BC00000
	jb	.LBB3_27
.LBB3_28:                               # %.loopexit
	testq	%r14, %r14
	je	.LBB3_31
# %bb.29:
	movl	$.L.str.26, %esi
	movq	%r14, %rdi
	callq	fopen
	movq	%rax, %r14
	testq	%rax, %rax
	jne	.LBB3_32
# %bb.30:
	movq	stderr(%rip), %rcx
	movl	$.L.str.27, %edi
	movl	$30, %esi
	jmp	.LBB3_23
.LBB3_31:
	movq	stdout(%rip), %r14
.LBB3_32:
	testl	%ebx, %ebx
	movl	(%rsp), %ebp                    # 4-byte Reload
	jle	.LBB3_46
# %bb.33:                               # %.lr.ph
	movl	%ebx, %ebx
	movl	$3435973837, %eax               # imm = 0xCCCCCCCD
	imulq	%rbx, %rax
	shrq	$35, %rax
	movq	%rax, 8(%rsp)                   # 8-byte Spill
	xorl	%r13d, %r13d
.LBB3_34:                               # =>This Inner Loop Header: Depth=1
	movq	(%r15,%r13,8), %r12
	testl	%ebp, %ebp
	je	.LBB3_35
# %bb.39:                               #   in Loop: Header=BB3_34 Depth=1
	cmpl	$1, %ebp
	jne	.LBB3_45
# %bb.40:                               #   in Loop: Header=BB3_34 Depth=1
	testq	%r12, %r12
	js	.LBB3_41
# %bb.42:                               #   in Loop: Header=BB3_34 Depth=1
	xorps	%xmm0, %xmm0
	cvtsi2ss	%r12, %xmm0
	jmp	.LBB3_43
.LBB3_35:                               #   in Loop: Header=BB3_34 Depth=1
	movl	%r13d, %eax
	xorl	%edx, %edx
	divl	8(%rsp)                         # 4-byte Folded Reload
	movq	stderr(%rip), %rax
	testq	%r13, %r13
	je	.LBB3_38
# %bb.36:                               #   in Loop: Header=BB3_34 Depth=1
	testl	%edx, %edx
	je	.LBB3_37
.LBB3_38:                               #   in Loop: Header=BB3_34 Depth=1
	movl	$46, %edi
	movq	%rax, %rsi
	callq	fputc@PLT
	jmp	.LBB3_44
.LBB3_41:                               #   in Loop: Header=BB3_34 Depth=1
	movq	%r12, %rax
	shrq	%rax
	movl	%r12d, %ecx
	andl	$1, %ecx
	orq	%rax, %rcx
	xorps	%xmm0, %xmm0
	cvtsi2ss	%rcx, %xmm0
	addss	%xmm0, %xmm0
.LBB3_43:                               #   in Loop: Header=BB3_34 Depth=1
	mulss	.LCPI3_0(%rip), %xmm0
	cvtss2sd	%xmm0, %xmm0
	movl	$.L.str.30, %edi
	movl	%r13d, %esi
	movq	%r12, %rdx
	movb	$1, %al
	callq	printf
	movl	$.L.str.31, %edi
	movq	%r12, %rsi
	xorl	%eax, %eax
	callq	printf
.LBB3_44:                               # %.thread103
                                        #   in Loop: Header=BB3_34 Depth=1
	movl	(%rsp), %ebp                    # 4-byte Reload
.LBB3_45:                               # %.thread103
                                        #   in Loop: Header=BB3_34 Depth=1
	leaq	20(%rsp), %rdi
	movq	%r12, %rsi
	movl	4(%rsp), %edx                   # 4-byte Reload
	movl	%ebp, %ecx
	callq	_Z15profileTransferm4Mode6Format
	movss	28(%rsp), %xmm2                 # xmm2 = mem[0],zero,zero,zero
	movss	32(%rsp), %xmm3                 # xmm3 = mem[0],zero,zero,zero
	xorps	%xmm0, %xmm0
	cvtss2sd	%xmm2, %xmm0
	xorps	%xmm1, %xmm1
	cvtss2sd	%xmm3, %xmm1
	addss	%xmm3, %xmm2
	cvtss2sd	%xmm2, %xmm2
	mulsd	.LCPI3_1(%rip), %xmm2
	movl	$.L.str.32, %esi
	movq	%r14, %rdi
	movq	%r12, %rdx
	movb	$3, %al
	callq	fprintf
	incq	%r13
	cmpq	%r13, %rbx
	jne	.LBB3_34
	jmp	.LBB3_46
.LBB3_37:                               #   in Loop: Header=BB3_34 Depth=1
	movl	$.L.str.28, %esi
	movq	%rax, %rdi
	movl	%r13d, %edx
	movq	%r12, %rcx
	xorl	%eax, %eax
	callq	fprintf
	jmp	.LBB3_44
.LBB3_46:                               # %._crit_edge
	testq	%r14, %r14
	je	.LBB3_48
# %bb.47:
	movq	%r14, %rdi
	callq	fclose
.LBB3_48:
	movq	%r15, %rdi
	callq	free
	xorl	%eax, %eax
	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
.LBB3_4:
	.cfi_def_cfa_offset 96
	movl	$.Lstr.2, %edi
	callq	puts@PLT
	movl	$-1, %edi
	callq	exit
.LBB3_16:
	callq	_Z9showUsagev
	xorl	%edi, %edi
	callq	exit
.LBB3_22:
	movq	stderr(%rip), %rcx
	movl	$.L.str.24, %edi
	movl	$53, %esi
.LBB3_23:
	movl	$1, %edx
	callq	fwrite@PLT
	movl	$1, %edi
	callq	exit
.Lfunc_end3:
	.size	main, .Lfunc_end3-main
	.cfi_endproc
	.section	.rodata,"a",@progbits
	.p2align	3, 0x0
.LJTI3_0:
	.quad	.LBB3_8
	.quad	.LBB3_4
	.quad	.LBB3_16
	.quad	.LBB3_4
	.quad	.LBB3_4
	.quad	.LBB3_4
	.quad	.LBB3_4
	.quad	.LBB3_5
	.quad	.LBB3_17
	.quad	.LBB3_14
                                        # -- End function
	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"num"
	.size	.L.str, 4

	.type	.L.str.1,@object                # @.str.1
.L.str.1:
	.asciz	"mode"
	.size	.L.str.1, 5

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	"format"
	.size	.L.str.2, 7

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

	.type	long_options,@object            # @long_options
	.data
	.globl	long_options
	.p2align	4, 0x0
long_options:
	.quad	.L.str
	.long	1                               # 0x1
	.zero	4
	.quad	0
	.long	0                               # 0x0
	.zero	4
	.quad	.L.str.1
	.long	1                               # 0x1
	.zero	4
	.quad	0
	.long	0                               # 0x0
	.zero	4
	.quad	.L.str.2
	.long	1                               # 0x1
	.zero	4
	.quad	0
	.long	0                               # 0x0
	.zero	4
	.quad	.L.str.3
	.long	0                               # 0x0
	.zero	4
	.quad	0
	.long	0                               # 0x0
	.zero	4
	.zero	32
	.size	long_options, 160

	.type	.L.str.4,@object                # @.str.4
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str.4:
	.asciz	"! host pinned memory allocation failed.\n! %s\n"
	.size	.L.str.4, 46

	.type	.L.str.6,@object                # @.str.6
.L.str.6:
	.asciz	" > allocating %s%d%s bytes in device memory\n"
	.size	.L.str.6, 45

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

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

	.type	.L.str.9,@object                # @.str.9
.L.str.9:
	.asciz	"! cuda error: %s\n"
	.size	.L.str.9, 18

	.type	.L.str.10,@object               # @.str.10
.L.str.10:
	.asciz	" > allocated in %.2f milliseconds\n"
	.size	.L.str.10, 35

	.type	.L.str.11,@object               # @.str.11
.L.str.11:
	.asciz	"  > copied host->device in %.2f milliseconds\n"
	.size	.L.str.11, 46

	.type	.L.str.12,@object               # @.str.12
.L.str.12:
	.asciz	"  > this equals a bandwidth of %s%.2f%s GB/s\n"
	.size	.L.str.12, 46

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

	.type	.L.str.14,@object               # @.str.14
.L.str.14:
	.asciz	"  > copied device->host in %.2f milliseconds\n"
	.size	.L.str.14, 46

	.type	.L.str.15,@object               # @.str.15
.L.str.15:
	.asciz	"  > this equals a bandwidth of %.2f GB/s\n"
	.size	.L.str.15, 42

	.type	.L.str.17,@object               # @.str.17
.L.str.17:
	.asciz	"n:m:f:ho:"
	.size	.L.str.17, 10

	.type	.L.str.18,@object               # @.str.18
.L.str.18:
	.asciz	"pinned"
	.size	.L.str.18, 7

	.type	.L.str.19,@object               # @.str.19
.L.str.19:
	.asciz	"pageable"
	.size	.L.str.19, 9

	.type	.L.str.20,@object               # @.str.20
.L.str.20:
	.asciz	"plot"
	.size	.L.str.20, 5

	.type	.L.str.21,@object               # @.str.21
.L.str.21:
	.asciz	"human"
	.size	.L.str.21, 6

	.type	.L.str.22,@object               # @.str.22
.L.str.22:
	.asciz	"plain"
	.size	.L.str.22, 6

	.type	.L.str.24,@object               # @.str.24
.L.str.24:
	.asciz	"!! can't provide amount of samples for human format.\n"
	.size	.L.str.24, 54

	.type	.L__const.main.samples_tmp,@object # @__const.main.samples_tmp
	.section	.rodata,"a",@progbits
	.p2align	4, 0x0
.L__const.main.samples_tmp:
	.quad	1                               # 0x1
	.quad	64                              # 0x40
	.quad	512                             # 0x200
	.quad	1024                            # 0x400
	.quad	1048576                         # 0x100000
	.quad	10485760                        # 0xa00000
	.quad	10485760                        # 0xa00000
	.quad	209715200                       # 0xc800000
	.quad	524288000                       # 0x1f400000
	.quad	943718400                       # 0x38400000
	.size	.L__const.main.samples_tmp, 80

	.type	.L.str.25,@object               # @.str.25
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str.25:
	.asciz	"num: %d\n"
	.size	.L.str.25, 9

	.type	.L.str.26,@object               # @.str.26
.L.str.26:
	.asciz	"a+"
	.size	.L.str.26, 3

	.type	.L.str.27,@object               # @.str.27
.L.str.27:
	.asciz	"!! error opening output-file.\n"
	.size	.L.str.27, 31

	.type	.L.str.28,@object               # @.str.28
.L.str.28:
	.asciz	"%d[%d]\n"
	.size	.L.str.28, 8

	.type	.L.str.30,@object               # @.str.30
.L.str.30:
	.asciz	"* test #%d: %d bytes ( %.2f MB)"
	.size	.L.str.30, 32

	.type	.L.str.31,@object               # @.str.31
.L.str.31:
	.asciz	"> allocating %d bytes in host memory\n"
	.size	.L.str.31, 38

	.type	.L.str.32,@object               # @.str.32
.L.str.32:
	.asciz	"%d\t%f\t%f\t%f\n"
	.size	.L.str.32, 13

	.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	"! host memory allocation failed!"
	.size	.Lstr, 33

	.type	.Lstr.1,@object                 # @str.1
.Lstr.1:
	.asciz	"usage: bandwidth [-p] [-o]"
	.size	.Lstr.1, 27

	.type	.Lstr.2,@object                 # @str.2
.Lstr.2:
	.asciz	"error"
	.size	.Lstr.2, 6

	.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 long_options
	.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
