// CopyBackAndForth.cu

#include <assert.h>
#include <stdio.h>

__device__ char devarray[16];

extern "C" void runTest()
{
   char zerobuf[16];
   memset(zerobuf, '@', sizeof(zerobuf));

   cudaError_t r = cudaMemcpyToSymbol(devarray, zerobuf, sizeof(zerobuf), 0, cudaMemcpyHostToDevice);
   assert(cudaSuccess == r);

   char out[16];
   r = cudaMemcpyFromSymbol(out, devarray, sizeof(devarray), 0, cudaMemcpyDeviceToHost);
   assert(cudaSuccess == r);

   assert(memcmp(out, zerobuf, sizeof(out)) == 0);
}

int main()
{
  runTest();
}

	.file	"tmpxft_0027c785_00000000-6_CopyBackAndForth.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
	.globl	runTest
	.type	runTest, @function
runTest:
.LFB2027:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	leaq	_ZL8devarray(%rip), %rbx
	xorl	%ecx, %ecx
	movl	$1, %r8d
	movl	$16, %edx
	movq	%rbx, %rdi
	subq	$48, %rsp
	.cfi_def_cfa_offset 64
	movaps	.LC0(%rip), %xmm0
	movq	%fs:40, %rax
	movq	%rax, 40(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rsi
	movups	%xmm0, 8(%rsp)
	call	cudaMemcpyToSymbol@PLT
	xorl	%ecx, %ecx
	leaq	24(%rsp), %rdi
	movq	%rbx, %rsi
	movl	$2, %r8d
	movl	$16, %edx
	call	cudaMemcpyFromSymbol@PLT
	movq	40(%rsp), %rax
	subq	%fs:40, %rax
	je	.L3
	call	__stack_chk_fail@PLT
.L3:
	addq	$48, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2027:
	.size	runTest, .-runTest
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB2028:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	call	runTest
	xorl	%eax, %eax
	popq	%rdx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2028:
	.size	main, .-main
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC1:
	.string	"devarray"
	.section	.text.startup
	.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
	pushq	$0
	.cfi_def_cfa_offset 24
	movl	$16, %r9d
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC1(%rip), %rdx
	movq	%rax, %rdi
	leaq	_ZL8devarray(%rip), %rsi
	movq	%rdx, %rcx
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterVar@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	popq	%rdx
	.cfi_def_cfa_offset 24
	popq	%rcx
	.cfi_def_cfa_offset 16
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rsi
	.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
	.local	_ZL8devarray
	.comm	_ZL8devarray,16,16
	.section	.rodata.cst16,"aM",@progbits,16
	.align 16
.LC0:
	.byte	64
	.byte	64
	.byte	64
	.byte	64
	.byte	64
	.byte	64
	.byte	64
	.byte	64
	.byte	64
	.byte	64
	.byte	64
	.byte	64
	.byte	64
	.byte	64
	.byte	64
	.byte	64
	.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:

// CopyBackAndForth.cu


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

__device__ char devarray[16];

extern "C" void runTest()
{
   char zerobuf[16];
   memset(zerobuf, '@', sizeof(zerobuf));

   hipError_t r = hipMemcpyToSymbol(HIP_SYMBOL(devarray), zerobuf, sizeof(zerobuf), 0, hipMemcpyHostToDevice);
   assert(hipSuccess == r);

   char out[16];
   r = hipMemcpyFromSymbol(out, HIP_SYMBOL(devarray), sizeof(devarray), 0, hipMemcpyDeviceToHost);
   assert(hipSuccess == r);

   assert(memcmp(out, zerobuf, sizeof(out)) == 0);
}

int main()
{
  runTest();
}

	.text
	.file	"CopyBackAndForth.hip"
	.section	.rodata.cst16,"aM",@progbits,16
	.p2align	4, 0x0                          # -- Begin function runTest
.LCPI0_0:
	.zero	16,64
	.text
	.globl	runTest
	.type	runTest,@function
runTest:                                # @runTest
	.cfi_startproc
# %bb.0:
	subq	$40, %rsp
	.cfi_def_cfa_offset 48
	movaps	.LCPI0_0(%rip), %xmm0           # xmm0 = [64,64,64,64,64,64,64,64,64,64,64,64,64,64,64,64]
	leaq	16(%rsp), %rsi
	movaps	%xmm0, (%rsi)
	movl	$devarray, %edi
	movl	$16, %edx
	xorl	%ecx, %ecx
	movl	$1, %r8d
	callq	hipMemcpyToSymbol
	movq	%rsp, %rdi
	movl	$devarray, %esi
	movl	$16, %edx
	xorl	%ecx, %ecx
	movl	$2, %r8d
	callq	hipMemcpyFromSymbol
	addq	$40, %rsp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end0:
	.size	runTest, .Lfunc_end0-runTest
	.cfi_endproc
                                        # -- End function
	.globl	main                            # -- Begin function main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%rax
	.cfi_def_cfa_offset 16
	callq	runTest
	xorl	%eax, %eax
	popq	%rcx
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end1:
	.size	main, .Lfunc_end1-main
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_ctor
	.type	__hip_module_ctor,@function
__hip_module_ctor:                      # @__hip_module_ctor
	.cfi_startproc
# %bb.0:
	pushq	%rax
	.cfi_def_cfa_offset 16
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	jne	.LBB2_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rdi
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB2_2:
	movl	$devarray, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movl	$16, %r9d
	xorl	%r8d, %r8d
	pushq	$0
	.cfi_adjust_cfa_offset 8
	pushq	$0
	.cfi_adjust_cfa_offset 8
	callq	__hipRegisterVar
	addq	$16, %rsp
	.cfi_adjust_cfa_offset -16
	movl	$__hip_module_dtor, %edi
	popq	%rax
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end2:
	.size	__hip_module_ctor, .Lfunc_end2-__hip_module_ctor
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_dtor
	.type	__hip_module_dtor,@function
__hip_module_dtor:                      # @__hip_module_dtor
	.cfi_startproc
# %bb.0:
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	je	.LBB3_2
# %bb.1:
	pushq	%rax
	.cfi_def_cfa_offset 16
	callq	__hipUnregisterFatBinary
	movq	$0, __hip_gpubin_handle(%rip)
	addq	$8, %rsp
	.cfi_def_cfa_offset 8
.LBB3_2:
	retq
.Lfunc_end3:
	.size	__hip_module_dtor, .Lfunc_end3-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	devarray,@object                # @devarray
	.local	devarray
	.comm	devarray,16,16
	.type	.L__unnamed_1,@object           # @0
	.section	.rodata.str1.1,"aMS",@progbits,1
.L__unnamed_1:
	.asciz	"devarray"
	.size	.L__unnamed_1, 9

	.type	__hip_fatbin_wrapper,@object    # @__hip_fatbin_wrapper
	.section	.hipFatBinSegment,"a",@progbits
	.p2align	3, 0x0
__hip_fatbin_wrapper:
	.long	1212764230                      # 0x48495046
	.long	1                               # 0x1
	.quad	__hip_fatbin
	.quad	0
	.size	__hip_fatbin_wrapper, 24

	.type	__hip_gpubin_handle,@object     # @__hip_gpubin_handle
	.local	__hip_gpubin_handle
	.comm	__hip_gpubin_handle,8,8
	.section	.init_array,"aw",@init_array
	.p2align	3, 0x0
	.quad	__hip_module_ctor
	.type	__hip_cuid_,@object             # @__hip_cuid_
	.bss
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               # 0x0
	.size	__hip_cuid_, 1

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

	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.protected	devarray                ; @devarray
	.type	devarray,@object
	.section	.bss,"aw",@nobits
	.globl	devarray
	.p2align	4, 0x0
devarray:
	.zero	16
	.size	devarray, 16

	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.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 devarray
	.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
