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

// This will output the proper CUDA error strings in the event that a CUDA host call returns an error
#define checkCudaErrors(err)           __checkCudaErrors (err, __FILE__, __LINE__)

inline void __checkCudaErrors(cudaError err, const char *file, const int line)
{
    if (cudaSuccess != err)
    {
        fprintf(stderr, "%s(%i) : CUDA Runtime API error %d: %s.\n",
                file, line, (int)err, cudaGetErrorString(err));
        exit(EXIT_FAILURE);
    }
}

int main(int argc, char **argv)
{
    printf("[%s] - Starting...\n", argv[0]);
    int gpuid[2] = {1, 2}; // we want to find the first two GPU's that can support P2P
    float total_time = 0.0;
    bool enable_p2p = true;

    if (enable_p2p)
    {
        // Enable peer access
        printf("Enabling peer access between GPU%d and GPU%d...\n", gpuid[0], gpuid[1]);
        checkCudaErrors(cudaSetDevice(gpuid[0]));
        checkCudaErrors(cudaDeviceEnablePeerAccess(gpuid[1], 0));
        checkCudaErrors(cudaSetDevice(gpuid[1]));
        checkCudaErrors(cudaDeviceEnablePeerAccess(gpuid[0], 0));
    }

    for (int loop = 0; loop < 100; loop++)
    {
        // Allocate buffers
        const size_t buf_size = 1024 * 1024 * 16 * sizeof(float);
        printf("Allocating buffers (%iMB on GPU%d, GPU%d and CPU Host)...\n", int(buf_size / 1024 / 1024), gpuid[0], gpuid[1]);
        checkCudaErrors(cudaSetDevice(gpuid[0]));
        float *g0;
        checkCudaErrors(cudaMalloc(&g0, buf_size));
        checkCudaErrors(cudaSetDevice(gpuid[1]));
        float *g1;
        checkCudaErrors(cudaMalloc(&g1, buf_size));

        // Create CUDA event handles
        cudaEvent_t start_event, stop_event;
        float time_memcpy;
        int eventflags = cudaEventBlockingSync;
        checkCudaErrors(cudaEventCreateWithFlags(&start_event, eventflags));
        checkCudaErrors(cudaEventCreateWithFlags(&stop_event, eventflags));

        // P2P memcopy() benchmark
        checkCudaErrors(cudaEventRecord(start_event, 0));

        for (int i = 0; i < 100; i++)
        {
            // With UVA we don't need to specify source and target devices, the
            // runtime figures this out by itself from the pointers

            // Ping-pong copy between GPUs
            if (i % 2 == 0)
            {
                checkCudaErrors(cudaMemcpy(g1, g0, buf_size, cudaMemcpyDefault));
            }
            else
            {
                checkCudaErrors(cudaMemcpy(g0, g1, buf_size, cudaMemcpyDefault));
            }
        }

        checkCudaErrors(cudaEventRecord(stop_event, 0));
        checkCudaErrors(cudaEventSynchronize(stop_event));
        checkCudaErrors(cudaEventElapsedTime(&time_memcpy, start_event, stop_event));
        total_time += time_memcpy;
        printf("cudaMemcpyPeer / cudaMemcpy (%f ms) between GPU%d and GPU%d: %.2fGB/s\n", time_memcpy, gpuid[0], gpuid[1],
               (1.0f / (time_memcpy / 1000.0f)) * ((100.0f * buf_size)) / 1024.0f / 1024.0f / 1024.0f);

        // Free resources
        checkCudaErrors(cudaEventDestroy(start_event));
        checkCudaErrors(cudaEventDestroy(stop_event));
        checkCudaErrors(cudaSetDevice(gpuid[0]));
        checkCudaErrors(cudaFree(g0));
        checkCudaErrors(cudaSetDevice(gpuid[1]));
        checkCudaErrors(cudaFree(g1));
    }


    if (enable_p2p)
    {
        // Disable peer access (also unregisters memory for non-UVA cases)
        printf("Disabling peer access...\n");
        checkCudaErrors(cudaSetDevice(gpuid[0]));
        checkCudaErrors(cudaDeviceDisablePeerAccess(gpuid[1]));
        checkCudaErrors(cudaSetDevice(gpuid[1]));
        checkCudaErrors(cudaDeviceDisablePeerAccess(gpuid[0]));
    }

    printf("Total time is %.2fs\n", total_time / 1000);
    //delete device_handler;
    return (EXIT_SUCCESS);
}

	.file	"tmpxft_00364b48_00000000-6_testP2P.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2031:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2031:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.section	.rodata._Z17__checkCudaErrors9cudaErrorPKci.str1.1,"aMS",@progbits,1
.LC0:
	.string	"%s(%i) : CUDA Runtime API error %d: %s.\n"
	.section	.text._Z17__checkCudaErrors9cudaErrorPKci,"axG",@progbits,_Z17__checkCudaErrors9cudaErrorPKci,comdat
	.weak	_Z17__checkCudaErrors9cudaErrorPKci
	.type	_Z17__checkCudaErrors9cudaErrorPKci, @function
_Z17__checkCudaErrors9cudaErrorPKci:
.LFB2027:
	.cfi_startproc
	endbr64
	testl	%edi, %edi
	je	.L2
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	movq	%rsi, %rbx
	subq	$16, %rsp
	.cfi_def_cfa_offset 32
	movl	%edx, 12(%rsp)
	movl	%edi, 8(%rsp)
	call	cudaGetErrorString@PLT
	movq	stderr(%rip), %rdi
	movq	%rbx, %rcx
	movl	$2, %esi
	pushq	%rdx
	.cfi_def_cfa_offset 40
	leaq	.LC0(%rip), %rdx
	pushq	%rax
	.cfi_def_cfa_offset 48
	movl	24(%rsp), %r9d
	xorl	%eax, %eax
	movl	28(%rsp), %r8d
	call	__fprintf_chk@PLT
	movl	$1, %edi
	call	exit@PLT
.L2:
	.cfi_def_cfa_offset 8
	.cfi_restore 3
	ret
	.cfi_endproc
.LFE2027:
	.size	_Z17__checkCudaErrors9cudaErrorPKci, .-_Z17__checkCudaErrors9cudaErrorPKci
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC2:
	.string	"[%s] - Starting...\n"
.LC3:
	.string	"Enabling peer access between GPU%d and GPU%d...\n"
.LC4:
	.string	"/home/ubuntu/Datasets/stackv2/train-structured/NanXiao/code-for-my-blog/master/2018/06/testP2P/testP2P.cu"
.LC5:
	.string	"Allocating buffers (%iMB on GPU%d, GPU%d and CPU Host)...\n"
.LC10:
	.string	"cudaMemcpyPeer / cudaMemcpy (%f ms) between GPU%d and GPU%d: %.2fGB/s\n"
.LC11:
	.string	"Disabling peer access...\n"
.LC12:
	.string	"Total time is %.2fs\n"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2028:
	.cfi_startproc
	endbr64
	pushq	%r14
	.cfi_def_cfa_offset 16
	.cfi_offset 14, -16
	movl	$2, %edi
	xorl	%r14d, %r14d
	pushq	%r13
	.cfi_def_cfa_offset 24
	.cfi_offset 13, -24
	pushq	%r12
	.cfi_def_cfa_offset 32
	.cfi_offset 12, -32
	leaq	.LC5(%rip), %r12
	pushq	%rbp
	.cfi_def_cfa_offset 40
	.cfi_offset 6, -40
	movl	$100, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 48
	.cfi_offset 3, -48
	leaq	.LC4(%rip), %rbx
	subq	$48, %rsp
	.cfi_def_cfa_offset 96
	movq	(%rsi), %rdx
	leaq	.LC2(%rip), %rsi
	movq	%fs:40, %rax
	movq	%rax, 40(%rsp)
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	$2, %ecx
	movl	$1, %edx
	xorl	%eax, %eax
	leaq	.LC3(%rip), %rsi
	movl	$2, %edi
	call	__printf_chk@PLT
	movl	$1, %edi
	call	cudaSetDevice@PLT
	movl	$28, %edx
	movq	%rbx, %rsi
	movl	%eax, %edi
	call	_Z17__checkCudaErrors9cudaErrorPKci
	xorl	%esi, %esi
	movl	$2, %edi
	call	cudaDeviceEnablePeerAccess@PLT
	movl	$29, %edx
	movq	%rbx, %rsi
	movl	%eax, %edi
	call	_Z17__checkCudaErrors9cudaErrorPKci
	movl	$2, %edi
	call	cudaSetDevice@PLT
	movl	$30, %edx
	movq	%rbx, %rsi
	movl	%eax, %edi
	call	_Z17__checkCudaErrors9cudaErrorPKci
	xorl	%esi, %esi
	movl	$1, %edi
	call	cudaDeviceEnablePeerAccess@PLT
	movl	$31, %edx
	movq	%rbx, %rsi
	movl	%eax, %edi
	call	_Z17__checkCudaErrors9cudaErrorPKci
.L12:
	movl	$2, %r8d
	movl	$1, %ecx
	movq	%r12, %rsi
	xorl	%eax, %eax
	movl	$64, %edx
	movl	$2, %edi
	xorl	%r13d, %r13d
	call	__printf_chk@PLT
	movl	$1, %edi
	call	cudaSetDevice@PLT
	movl	$39, %edx
	movq	%rbx, %rsi
	movl	%eax, %edi
	call	_Z17__checkCudaErrors9cudaErrorPKci
	leaq	8(%rsp), %rdi
	movl	$67108864, %esi
	call	cudaMalloc@PLT
	movl	$41, %edx
	movq	%rbx, %rsi
	movl	%eax, %edi
	call	_Z17__checkCudaErrors9cudaErrorPKci
	movl	$2, %edi
	call	cudaSetDevice@PLT
	movl	$42, %edx
	movq	%rbx, %rsi
	movl	%eax, %edi
	call	_Z17__checkCudaErrors9cudaErrorPKci
	leaq	16(%rsp), %rdi
	movl	$67108864, %esi
	call	cudaMalloc@PLT
	movl	$44, %edx
	movq	%rbx, %rsi
	movl	%eax, %edi
	call	_Z17__checkCudaErrors9cudaErrorPKci
	leaq	24(%rsp), %rdi
	movl	$1, %esi
	call	cudaEventCreateWithFlags@PLT
	movl	$50, %edx
	movq	%rbx, %rsi
	movl	%eax, %edi
	call	_Z17__checkCudaErrors9cudaErrorPKci
	leaq	32(%rsp), %rdi
	movl	$1, %esi
	call	cudaEventCreateWithFlags@PLT
	movl	$51, %edx
	movq	%rbx, %rsi
	movl	%eax, %edi
	call	_Z17__checkCudaErrors9cudaErrorPKci
	movq	24(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movl	$54, %edx
	movq	%rbx, %rsi
	movl	%eax, %edi
	call	_Z17__checkCudaErrors9cudaErrorPKci
.L11:
	movq	16(%rsp), %rax
	movq	8(%rsp), %rdi
	movl	$4, %ecx
	movl	$67108864, %edx
	testl	$1, %r13d
	jne	.L9
	movq	%rdi, %rsi
	movq	%rax, %rdi
	call	cudaMemcpy@PLT
	movl	$64, %edx
	movl	%eax, %edi
	jmp	.L17
.L9:
	movq	%rax, %rsi
	call	cudaMemcpy@PLT
	movl	$68, %edx
	movl	%eax, %edi
.L17:
	movq	%rbx, %rsi
	incl	%r13d
	call	_Z17__checkCudaErrors9cudaErrorPKci
	cmpl	$100, %r13d
	jne	.L11
	movq	32(%rsp), %rdi
	xorl	%esi, %esi
	call	cudaEventRecord@PLT
	movl	$72, %edx
	movq	%rbx, %rsi
	movl	%eax, %edi
	call	_Z17__checkCudaErrors9cudaErrorPKci
	movq	32(%rsp), %rdi
	call	cudaEventSynchronize@PLT
	movl	$73, %edx
	movq	%rbx, %rsi
	movl	%eax, %edi
	call	_Z17__checkCudaErrors9cudaErrorPKci
	movq	32(%rsp), %rdx
	movq	24(%rsp), %rsi
	leaq	4(%rsp), %rdi
	call	cudaEventElapsedTime@PLT
	movl	$74, %edx
	movq	%rbx, %rsi
	movl	%eax, %edi
	call	_Z17__checkCudaErrors9cudaErrorPKci
	movss	4(%rsp), %xmm0
	movd	%r14d, %xmm3
	movss	.LC7(%rip), %xmm1
	movl	$2, %ecx
	movl	$1, %edx
	leaq	.LC10(%rip), %rsi
	movb	$2, %al
	movaps	%xmm0, %xmm2
	divss	.LC6(%rip), %xmm2
	addss	%xmm0, %xmm3
	movl	$2, %edi
	divss	%xmm2, %xmm1
	movss	.LC9(%rip), %xmm2
	cvtss2sd	%xmm0, %xmm0
	movd	%xmm3, %r14d
	mulss	.LC8(%rip), %xmm1
	mulss	%xmm2, %xmm1
	mulss	%xmm2, %xmm1
	mulss	%xmm2, %xmm1
	cvtss2sd	%xmm1, %xmm1
	call	__printf_chk@PLT
	movq	24(%rsp), %rdi
	call	cudaEventDestroy@PLT
	movl	$80, %edx
	movq	%rbx, %rsi
	movl	%eax, %edi
	call	_Z17__checkCudaErrors9cudaErrorPKci
	movq	32(%rsp), %rdi
	call	cudaEventDestroy@PLT
	movl	$81, %edx
	movq	%rbx, %rsi
	movl	%eax, %edi
	call	_Z17__checkCudaErrors9cudaErrorPKci
	movl	$1, %edi
	call	cudaSetDevice@PLT
	movl	$82, %edx
	movq	%rbx, %rsi
	movl	%eax, %edi
	call	_Z17__checkCudaErrors9cudaErrorPKci
	movq	8(%rsp), %rdi
	call	cudaFree@PLT
	movl	$83, %edx
	movq	%rbx, %rsi
	movl	%eax, %edi
	call	_Z17__checkCudaErrors9cudaErrorPKci
	movl	$2, %edi
	call	cudaSetDevice@PLT
	movl	$84, %edx
	movq	%rbx, %rsi
	movl	%eax, %edi
	call	_Z17__checkCudaErrors9cudaErrorPKci
	movq	16(%rsp), %rdi
	call	cudaFree@PLT
	movl	$85, %edx
	movq	%rbx, %rsi
	movl	%eax, %edi
	call	_Z17__checkCudaErrors9cudaErrorPKci
	decl	%ebp
	jne	.L12
	leaq	.LC11(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	$1, %edi
	call	cudaSetDevice@PLT
	movl	$93, %edx
	movq	%rbx, %rsi
	movl	%eax, %edi
	call	_Z17__checkCudaErrors9cudaErrorPKci
	movl	$2, %edi
	call	cudaDeviceDisablePeerAccess@PLT
	movl	$94, %edx
	movq	%rbx, %rsi
	movl	%eax, %edi
	call	_Z17__checkCudaErrors9cudaErrorPKci
	movl	$2, %edi
	call	cudaSetDevice@PLT
	movl	$95, %edx
	movq	%rbx, %rsi
	movl	%eax, %edi
	call	_Z17__checkCudaErrors9cudaErrorPKci
	movl	$1, %edi
	call	cudaDeviceDisablePeerAccess@PLT
	movq	%rbx, %rsi
	movl	$96, %edx
	movl	%eax, %edi
	call	_Z17__checkCudaErrors9cudaErrorPKci
	leaq	.LC12(%rip), %rsi
	movl	$2, %edi
	movb	$1, %al
	movd	%r14d, %xmm0
	divss	.LC6(%rip), %xmm0
	cvtss2sd	%xmm0, %xmm0
	call	__printf_chk@PLT
	movq	40(%rsp), %rax
	subq	%fs:40, %rax
	je	.L13
	call	__stack_chk_fail@PLT
.L13:
	addq	$48, %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
.LFE2028:
	.size	main, .-main
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2054:
	.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
.LFE2054:
	.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
.LC6:
	.long	1148846080
	.align 4
.LC7:
	.long	1065353216
	.align 4
.LC8:
	.long	1338507264
	.align 4
.LC9:
	.long	981467136
	.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 <hip/hip_runtime.h>
#include <stdlib.h>
#include <stdio.h>

// This will output the proper CUDA error strings in the event that a CUDA host call returns an error
#define checkCudaErrors(err)           __checkCudaErrors (err, __FILE__, __LINE__)

inline void __checkCudaErrors(hipError_t err, const char *file, const int line)
{
    if (hipSuccess != err)
    {
        fprintf(stderr, "%s(%i) : CUDA Runtime API error %d: %s.\n",
                file, line, (int)err, hipGetErrorString(err));
        exit(EXIT_FAILURE);
    }
}

int main(int argc, char **argv)
{
    printf("[%s] - Starting...\n", argv[0]);
    int gpuid[2] = {1, 2}; // we want to find the first two GPU's that can support P2P
    float total_time = 0.0;
    bool enable_p2p = true;

    if (enable_p2p)
    {
        // Enable peer access
        printf("Enabling peer access between GPU%d and GPU%d...\n", gpuid[0], gpuid[1]);
        checkCudaErrors(hipSetDevice(gpuid[0]));
        checkCudaErrors(hipDeviceEnablePeerAccess(gpuid[1], 0));
        checkCudaErrors(hipSetDevice(gpuid[1]));
        checkCudaErrors(hipDeviceEnablePeerAccess(gpuid[0], 0));
    }

    for (int loop = 0; loop < 100; loop++)
    {
        // Allocate buffers
        const size_t buf_size = 1024 * 1024 * 16 * sizeof(float);
        printf("Allocating buffers (%iMB on GPU%d, GPU%d and CPU Host)...\n", int(buf_size / 1024 / 1024), gpuid[0], gpuid[1]);
        checkCudaErrors(hipSetDevice(gpuid[0]));
        float *g0;
        checkCudaErrors(hipMalloc(&g0, buf_size));
        checkCudaErrors(hipSetDevice(gpuid[1]));
        float *g1;
        checkCudaErrors(hipMalloc(&g1, buf_size));

        // Create CUDA event handles
        hipEvent_t start_event, stop_event;
        float time_memcpy;
        int eventflags = hipEventBlockingSync;
        checkCudaErrors(hipEventCreateWithFlags(&start_event, eventflags));
        checkCudaErrors(hipEventCreateWithFlags(&stop_event, eventflags));

        // P2P memcopy() benchmark
        checkCudaErrors(hipEventRecord(start_event, 0));

        for (int i = 0; i < 100; i++)
        {
            // With UVA we don't need to specify source and target devices, the
            // runtime figures this out by itself from the pointers

            // Ping-pong copy between GPUs
            if (i % 2 == 0)
            {
                checkCudaErrors(hipMemcpy(g1, g0, buf_size, hipMemcpyDefault));
            }
            else
            {
                checkCudaErrors(hipMemcpy(g0, g1, buf_size, hipMemcpyDefault));
            }
        }

        checkCudaErrors(hipEventRecord(stop_event, 0));
        checkCudaErrors(hipEventSynchronize(stop_event));
        checkCudaErrors(hipEventElapsedTime(&time_memcpy, start_event, stop_event));
        total_time += time_memcpy;
        printf("hipMemcpyPeer / hipMemcpy (%f ms) between GPU%d and GPU%d: %.2fGB/s\n", time_memcpy, gpuid[0], gpuid[1],
               (1.0f / (time_memcpy / 1000.0f)) * ((100.0f * buf_size)) / 1024.0f / 1024.0f / 1024.0f);

        // Free resources
        checkCudaErrors(hipEventDestroy(start_event));
        checkCudaErrors(hipEventDestroy(stop_event));
        checkCudaErrors(hipSetDevice(gpuid[0]));
        checkCudaErrors(hipFree(g0));
        checkCudaErrors(hipSetDevice(gpuid[1]));
        checkCudaErrors(hipFree(g1));
    }


    if (enable_p2p)
    {
        // Disable peer access (also unregisters memory for non-UVA cases)
        printf("Disabling peer access...\n");
        checkCudaErrors(hipSetDevice(gpuid[0]));
        checkCudaErrors(hipDeviceDisablePeerAccess(gpuid[1]));
        checkCudaErrors(hipSetDevice(gpuid[1]));
        checkCudaErrors(hipDeviceDisablePeerAccess(gpuid[0]));
    }

    printf("Total time is %.2fs\n", total_time / 1000);
    //delete device_handler;
    return (EXIT_SUCCESS);
}

	.text
	.file	"testP2P.hip"
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0                          # -- Begin function main
.LCPI0_0:
	.long	0x447a0000                      # float 1000
.LCPI0_1:
	.long	0x3f800000                      # float 1
.LCPI0_2:
	.long	0x4fc80000                      # float 6.7108864E+9
.LCPI0_3:
	.long	0x3a800000                      # float 9.765625E-4
	.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), %rsi
	movl	$.L.str, %edi
	xorl	%eax, %eax
	callq	printf
	movl	$.L.str.1, %edi
	movl	$1, %esi
	movl	$2, %edx
	xorl	%eax, %eax
	callq	printf
	movl	$1, %edi
	callq	hipSetDevice
	testl	%eax, %eax
	jne	.LBB0_1
# %bb.4:                                # %_Z17__checkCudaErrors10hipError_tPKci.exit
	movl	$2, %edi
	xorl	%esi, %esi
	callq	hipDeviceEnablePeerAccess
	testl	%eax, %eax
	jne	.LBB0_5
# %bb.6:                                # %_Z17__checkCudaErrors10hipError_tPKci.exit31
	movl	$2, %edi
	callq	hipSetDevice
	testl	%eax, %eax
	jne	.LBB0_7
# %bb.8:                                # %_Z17__checkCudaErrors10hipError_tPKci.exit33
	movl	$1, %edi
	xorl	%esi, %esi
	callq	hipDeviceEnablePeerAccess
	testl	%eax, %eax
	jne	.LBB0_60
# %bb.9:                                # %_Z17__checkCudaErrors10hipError_tPKci.exit35.preheader.preheader
	xorps	%xmm0, %xmm0
	movss	%xmm0, 12(%rsp)                 # 4-byte Spill
	leaq	24(%rsp), %r15
	leaq	16(%rsp), %r12
	leaq	52(%rsp), %r13
	xorl	%ebx, %ebx
.LBB0_10:                               # %_Z17__checkCudaErrors10hipError_tPKci.exit35.preheader
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB0_25 Depth 2
	movl	$.L.str.3, %edi
	movl	$64, %esi
	movl	$1, %edx
	movl	$2, %ecx
	xorl	%eax, %eax
	callq	printf
	movl	$1, %edi
	callq	hipSetDevice
	testl	%eax, %eax
	jne	.LBB0_11
# %bb.13:                               # %_Z17__checkCudaErrors10hipError_tPKci.exit37
                                        #   in Loop: Header=BB0_10 Depth=1
	movl	$67108864, %esi                 # imm = 0x4000000
	leaq	40(%rsp), %rdi
	callq	hipMalloc
	testl	%eax, %eax
	jne	.LBB0_14
# %bb.15:                               # %_Z17__checkCudaErrors10hipError_tPKci.exit39
                                        #   in Loop: Header=BB0_10 Depth=1
	movl	$2, %edi
	callq	hipSetDevice
	testl	%eax, %eax
	jne	.LBB0_16
# %bb.17:                               # %_Z17__checkCudaErrors10hipError_tPKci.exit41
                                        #   in Loop: Header=BB0_10 Depth=1
	movl	$67108864, %esi                 # imm = 0x4000000
	leaq	32(%rsp), %rdi
	callq	hipMalloc
	testl	%eax, %eax
	jne	.LBB0_18
# %bb.19:                               # %_Z17__checkCudaErrors10hipError_tPKci.exit43
                                        #   in Loop: Header=BB0_10 Depth=1
	movq	%r15, %rdi
	movl	$1, %esi
	callq	hipEventCreateWithFlags
	testl	%eax, %eax
	jne	.LBB0_20
# %bb.21:                               # %_Z17__checkCudaErrors10hipError_tPKci.exit45
                                        #   in Loop: Header=BB0_10 Depth=1
	movq	%r12, %rdi
	movl	$1, %esi
	callq	hipEventCreateWithFlags
	testl	%eax, %eax
	jne	.LBB0_22
# %bb.23:                               # %_Z17__checkCudaErrors10hipError_tPKci.exit47
                                        #   in Loop: Header=BB0_10 Depth=1
	movq	24(%rsp), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	testl	%eax, %eax
	jne	.LBB0_28
# %bb.24:                               # %_Z17__checkCudaErrors10hipError_tPKci.exit49.preheader.preheader
                                        #   in Loop: Header=BB0_10 Depth=1
	xorl	%r14d, %r14d
.LBB0_25:                               # %_Z17__checkCudaErrors10hipError_tPKci.exit49.preheader
                                        #   Parent Loop BB0_10 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	testb	$1, %r14b
	jne	.LBB0_48
# %bb.26:                               #   in Loop: Header=BB0_25 Depth=2
	movq	32(%rsp), %rdi
	movq	40(%rsp), %rsi
	movl	$67108864, %edx                 # imm = 0x4000000
	movl	$4, %ecx
	callq	hipMemcpy
	testl	%eax, %eax
	je	.LBB0_29
	jmp	.LBB0_27
.LBB0_48:                               #   in Loop: Header=BB0_25 Depth=2
	movq	40(%rsp), %rdi
	movq	32(%rsp), %rsi
	movl	$67108864, %edx                 # imm = 0x4000000
	movl	$4, %ecx
	callq	hipMemcpy
	testl	%eax, %eax
	jne	.LBB0_49
.LBB0_29:                               # %_Z17__checkCudaErrors10hipError_tPKci.exit69
                                        #   in Loop: Header=BB0_25 Depth=2
	incl	%r14d
	cmpl	$100, %r14d
	jne	.LBB0_25
# %bb.30:                               #   in Loop: Header=BB0_10 Depth=1
	movq	16(%rsp), %rdi
	xorl	%esi, %esi
	callq	hipEventRecord
	testl	%eax, %eax
	jne	.LBB0_31
# %bb.32:                               # %_Z17__checkCudaErrors10hipError_tPKci.exit51
                                        #   in Loop: Header=BB0_10 Depth=1
	movq	16(%rsp), %rdi
	callq	hipEventSynchronize
	testl	%eax, %eax
	jne	.LBB0_33
# %bb.34:                               # %_Z17__checkCudaErrors10hipError_tPKci.exit53
                                        #   in Loop: Header=BB0_10 Depth=1
	movq	24(%rsp), %rsi
	movq	16(%rsp), %rdx
	movq	%r13, %rdi
	callq	hipEventElapsedTime
	testl	%eax, %eax
	jne	.LBB0_35
# %bb.36:                               # %_Z17__checkCudaErrors10hipError_tPKci.exit55
                                        #   in Loop: Header=BB0_10 Depth=1
	movss	52(%rsp), %xmm1                 # xmm1 = mem[0],zero,zero,zero
	cvtss2sd	%xmm1, %xmm0
	movss	%xmm1, 48(%rsp)                 # 4-byte Spill
	divss	.LCPI0_0(%rip), %xmm1
	movss	.LCPI0_1(%rip), %xmm2           # xmm2 = mem[0],zero,zero,zero
	divss	%xmm1, %xmm2
	mulss	.LCPI0_2(%rip), %xmm2
	movss	.LCPI0_3(%rip), %xmm1           # xmm1 = mem[0],zero,zero,zero
	mulss	%xmm1, %xmm2
	mulss	%xmm1, %xmm2
	mulss	%xmm1, %xmm2
	xorps	%xmm1, %xmm1
	cvtss2sd	%xmm2, %xmm1
	movl	$.L.str.4, %edi
	movl	$1, %esi
	movl	$2, %edx
	movb	$2, %al
	callq	printf
	movq	24(%rsp), %rdi
	callq	hipEventDestroy
	testl	%eax, %eax
	jne	.LBB0_37
# %bb.38:                               # %_Z17__checkCudaErrors10hipError_tPKci.exit57
                                        #   in Loop: Header=BB0_10 Depth=1
	movq	16(%rsp), %rdi
	callq	hipEventDestroy
	testl	%eax, %eax
	jne	.LBB0_39
# %bb.40:                               # %_Z17__checkCudaErrors10hipError_tPKci.exit59
                                        #   in Loop: Header=BB0_10 Depth=1
	movl	$1, %edi
	callq	hipSetDevice
	testl	%eax, %eax
	jne	.LBB0_41
# %bb.42:                               # %_Z17__checkCudaErrors10hipError_tPKci.exit61
                                        #   in Loop: Header=BB0_10 Depth=1
	movq	40(%rsp), %rdi
	callq	hipFree
	testl	%eax, %eax
	jne	.LBB0_43
# %bb.44:                               # %_Z17__checkCudaErrors10hipError_tPKci.exit63
                                        #   in Loop: Header=BB0_10 Depth=1
	movl	$2, %edi
	callq	hipSetDevice
	testl	%eax, %eax
	jne	.LBB0_45
# %bb.46:                               # %_Z17__checkCudaErrors10hipError_tPKci.exit65
                                        #   in Loop: Header=BB0_10 Depth=1
	movq	32(%rsp), %rdi
	callq	hipFree
	testl	%eax, %eax
	jne	.LBB0_47
# %bb.50:                               # %_Z17__checkCudaErrors10hipError_tPKci.exit67
                                        #   in Loop: Header=BB0_10 Depth=1
	movss	12(%rsp), %xmm0                 # 4-byte Reload
                                        # xmm0 = mem[0],zero,zero,zero
	addss	48(%rsp), %xmm0                 # 4-byte Folded Reload
	movss	%xmm0, 12(%rsp)                 # 4-byte Spill
	incl	%ebx
	cmpl	$100, %ebx
	jne	.LBB0_10
# %bb.51:
	movl	$.Lstr, %edi
	callq	puts@PLT
	movl	$1, %edi
	callq	hipSetDevice
	testl	%eax, %eax
	jne	.LBB0_52
# %bb.53:                               # %_Z17__checkCudaErrors10hipError_tPKci.exit73
	movl	$2, %edi
	callq	hipDeviceDisablePeerAccess
	testl	%eax, %eax
	jne	.LBB0_54
# %bb.55:                               # %_Z17__checkCudaErrors10hipError_tPKci.exit75
	movl	$2, %edi
	callq	hipSetDevice
	testl	%eax, %eax
	jne	.LBB0_56
# %bb.57:                               # %_Z17__checkCudaErrors10hipError_tPKci.exit77
	movl	$1, %edi
	callq	hipDeviceDisablePeerAccess
	testl	%eax, %eax
	jne	.LBB0_58
# %bb.59:                               # %_Z17__checkCudaErrors10hipError_tPKci.exit79
	movss	12(%rsp), %xmm0                 # 4-byte Reload
                                        # xmm0 = mem[0],zero,zero,zero
	divss	.LCPI0_0(%rip), %xmm0
	cvtss2sd	%xmm0, %xmm0
	movl	$.L.str.6, %edi
	movb	$1, %al
	callq	printf
	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
.LBB0_27:
	.cfi_def_cfa_offset 112
	movl	%eax, %ebp
	movq	stderr(%rip), %rbx
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.7, %esi
	movl	$.L.str.2, %edx
	movq	%rbx, %rdi
	movl	$66, %ecx
	jmp	.LBB0_12
.LBB0_49:
	movl	%eax, %ebp
	movq	stderr(%rip), %rbx
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.7, %esi
	movl	$.L.str.2, %edx
	movq	%rbx, %rdi
	movl	$70, %ecx
.LBB0_12:
	movl	%ebp, %r8d
.LBB0_3:
	movq	%rax, %r9
	xorl	%eax, %eax
	callq	fprintf
	movl	$1, %edi
	callq	exit
.LBB0_28:
	movl	%eax, %ebp
	movq	stderr(%rip), %rbx
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.7, %esi
	movl	$.L.str.2, %edx
	movq	%rbx, %rdi
	movl	$56, %ecx
	jmp	.LBB0_12
.LBB0_22:
	movl	%eax, %ebp
	movq	stderr(%rip), %rbx
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.7, %esi
	movl	$.L.str.2, %edx
	movq	%rbx, %rdi
	movl	$53, %ecx
	jmp	.LBB0_12
.LBB0_20:
	movl	%eax, %ebp
	movq	stderr(%rip), %rbx
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.7, %esi
	movl	$.L.str.2, %edx
	movq	%rbx, %rdi
	movl	$52, %ecx
	jmp	.LBB0_12
.LBB0_18:
	movl	%eax, %ebp
	movq	stderr(%rip), %rbx
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.7, %esi
	movl	$.L.str.2, %edx
	movq	%rbx, %rdi
	movl	$46, %ecx
	jmp	.LBB0_12
.LBB0_16:
	movl	%eax, %ebp
	movq	stderr(%rip), %rbx
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.7, %esi
	movl	$.L.str.2, %edx
	movq	%rbx, %rdi
	movl	$44, %ecx
	jmp	.LBB0_12
.LBB0_14:
	movl	%eax, %ebp
	movq	stderr(%rip), %rbx
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.7, %esi
	movl	$.L.str.2, %edx
	movq	%rbx, %rdi
	movl	$43, %ecx
	jmp	.LBB0_12
.LBB0_11:
	movl	%eax, %ebp
	movq	stderr(%rip), %rbx
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.7, %esi
	movl	$.L.str.2, %edx
	movq	%rbx, %rdi
	movl	$41, %ecx
	jmp	.LBB0_12
.LBB0_47:
	movl	%eax, %ebp
	movq	stderr(%rip), %rbx
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.7, %esi
	movl	$.L.str.2, %edx
	movq	%rbx, %rdi
	movl	$87, %ecx
	jmp	.LBB0_12
.LBB0_45:
	movl	%eax, %ebp
	movq	stderr(%rip), %rbx
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.7, %esi
	movl	$.L.str.2, %edx
	movq	%rbx, %rdi
	movl	$86, %ecx
	jmp	.LBB0_12
.LBB0_43:
	movl	%eax, %ebp
	movq	stderr(%rip), %rbx
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.7, %esi
	movl	$.L.str.2, %edx
	movq	%rbx, %rdi
	movl	$85, %ecx
	jmp	.LBB0_12
.LBB0_41:
	movl	%eax, %ebp
	movq	stderr(%rip), %rbx
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.7, %esi
	movl	$.L.str.2, %edx
	movq	%rbx, %rdi
	movl	$84, %ecx
	jmp	.LBB0_12
.LBB0_39:
	movl	%eax, %ebp
	movq	stderr(%rip), %rbx
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.7, %esi
	movl	$.L.str.2, %edx
	movq	%rbx, %rdi
	movl	$83, %ecx
	jmp	.LBB0_12
.LBB0_37:
	movl	%eax, %ebp
	movq	stderr(%rip), %rbx
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.7, %esi
	movl	$.L.str.2, %edx
	movq	%rbx, %rdi
	movl	$82, %ecx
	jmp	.LBB0_12
.LBB0_35:
	movl	%eax, %ebp
	movq	stderr(%rip), %rbx
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.7, %esi
	movl	$.L.str.2, %edx
	movq	%rbx, %rdi
	movl	$76, %ecx
	jmp	.LBB0_12
.LBB0_33:
	movl	%eax, %ebp
	movq	stderr(%rip), %rbx
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.7, %esi
	movl	$.L.str.2, %edx
	movq	%rbx, %rdi
	movl	$75, %ecx
	jmp	.LBB0_12
.LBB0_31:
	movl	%eax, %ebp
	movq	stderr(%rip), %rbx
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.7, %esi
	movl	$.L.str.2, %edx
	movq	%rbx, %rdi
	movl	$74, %ecx
	jmp	.LBB0_12
.LBB0_1:
	movl	%eax, %ebx
	movq	stderr(%rip), %r14
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.7, %esi
	movl	$.L.str.2, %edx
	movq	%r14, %rdi
	movl	$30, %ecx
	jmp	.LBB0_2
.LBB0_5:
	movl	%eax, %ebx
	movq	stderr(%rip), %r14
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.7, %esi
	movl	$.L.str.2, %edx
	movq	%r14, %rdi
	movl	$31, %ecx
	jmp	.LBB0_2
.LBB0_7:
	movl	%eax, %ebx
	movq	stderr(%rip), %r14
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.7, %esi
	movl	$.L.str.2, %edx
	movq	%r14, %rdi
	movl	$32, %ecx
	jmp	.LBB0_2
.LBB0_60:
	movl	%eax, %ebx
	movq	stderr(%rip), %r14
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.7, %esi
	movl	$.L.str.2, %edx
	movq	%r14, %rdi
	movl	$33, %ecx
	jmp	.LBB0_2
.LBB0_52:
	movl	%eax, %ebx
	movq	stderr(%rip), %r14
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.7, %esi
	movl	$.L.str.2, %edx
	movq	%r14, %rdi
	movl	$95, %ecx
	jmp	.LBB0_2
.LBB0_54:
	movl	%eax, %ebx
	movq	stderr(%rip), %r14
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.7, %esi
	movl	$.L.str.2, %edx
	movq	%r14, %rdi
	movl	$96, %ecx
	jmp	.LBB0_2
.LBB0_56:
	movl	%eax, %ebx
	movq	stderr(%rip), %r14
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.7, %esi
	movl	$.L.str.2, %edx
	movq	%r14, %rdi
	movl	$97, %ecx
	jmp	.LBB0_2
.LBB0_58:
	movl	%eax, %ebx
	movq	stderr(%rip), %r14
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.7, %esi
	movl	$.L.str.2, %edx
	movq	%r14, %rdi
	movl	$98, %ecx
.LBB0_2:
	movl	%ebx, %r8d
	jmp	.LBB0_3
.Lfunc_end0:
	.size	main, .Lfunc_end0-main
	.cfi_endproc
                                        # -- End function
	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"[%s] - Starting...\n"
	.size	.L.str, 20

	.type	.L.str.1,@object                # @.str.1
.L.str.1:
	.asciz	"Enabling peer access between GPU%d and GPU%d...\n"
	.size	.L.str.1, 49

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	"/home/ubuntu/Datasets/stackv2/train-structured-repos-hip-sm89/NanXiao/code-for-my-blog/master/2018/06/testP2P/testP2P.hip"
	.size	.L.str.2, 122

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"Allocating buffers (%iMB on GPU%d, GPU%d and CPU Host)...\n"
	.size	.L.str.3, 59

	.type	.L.str.4,@object                # @.str.4
.L.str.4:
	.asciz	"hipMemcpyPeer / hipMemcpy (%f ms) between GPU%d and GPU%d: %.2fGB/s\n"
	.size	.L.str.4, 69

	.type	.L.str.6,@object                # @.str.6
.L.str.6:
	.asciz	"Total time is %.2fs\n"
	.size	.L.str.6, 21

	.type	.L.str.7,@object                # @.str.7
.L.str.7:
	.asciz	"%s(%i) : CUDA Runtime API error %d: %s.\n"
	.size	.L.str.7, 41

	.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	"Disabling peer access..."
	.size	.Lstr, 25

	.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
