// What if we are given a device pointer that is offset from any of the device pointers we provided to the client?
//
// This file is a test-case for this.  Then we can look at handling that...

#include <iostream>
#include <memory>
#include <cassert>

using namespace std;

#include <cuda.h>

// __global__ void getValue(float *data, int idx, float value) {
//     if(threadIdx.x == 0) {
//         data[idx] += value;
//     }
// }

int main(int argc, char *argv[]) {
    int N = 1024;

    CUstream stream;
    cuStreamCreate(&stream, 0);

    float *hostFloats;
    cuMemHostAlloc((void **)&hostFloats, N * sizeof(float), CU_MEMHOSTALLOC_PORTABLE);

    CUdeviceptr deviceFloats;
    cuMemAlloc(&deviceFloats, N * sizeof(float));

    hostFloats[128] = 123.456f;
    hostFloats[129] = 444.0f;
    hostFloats[130] = 321.0f;
    hostFloats[131] = 111.0f;

    // now we will copy 16 bytes, starting at location 128...
    cuMemcpyHtoDAsync(
        (CUdeviceptr)(((float *)deviceFloats) + 64),
        hostFloats + 128,
        4 * sizeof(float),
        stream
    );
    cuStreamSynchronize(stream);
    // now copy back entire buffer
    hostFloats[64] = 0.0f;
    hostFloats[65] = 0.0f;
    hostFloats[66] = 0.0f;
    hostFloats[67] = 0.0f;
    cuMemcpyDtoHAsync(hostFloats, deviceFloats, N * sizeof(float), stream);
    cuStreamSynchronize(stream);

    // and check the values...
    cout << hostFloats[64] << endl;
    cout << hostFloats[65] << endl;
    cout << hostFloats[66] << endl;
    cout << hostFloats[67] << endl;

    assert(hostFloats[64] == 123.456f);
    assert(hostFloats[65] == 444.0f);
    assert(hostFloats[66] == 321);
    assert(hostFloats[67] == 111);

    cuMemFreeHost(hostFloats);
    cuMemFree(deviceFloats);
    cuStreamDestroy(stream);

    return 0;
}

	.file	"tmpxft_0022ccd8_00000000-6_testpartialcopy.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB4282:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE4282:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB4279:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	xorl	%esi, %esi
	leaq	_ZSt4cout(%rip), %rbx
	subq	$32, %rsp
	.cfi_def_cfa_offset 48
	movq	%fs:40, %rax
	movq	%rax, 24(%rsp)
	xorl	%eax, %eax
	movq	%rsp, %rdi
	call	cuStreamCreate@PLT
	movl	$1, %edx
	leaq	8(%rsp), %rdi
	movl	$4096, %esi
	call	cuMemHostAlloc@PLT
	leaq	16(%rsp), %rdi
	movl	$4096, %esi
	call	cuMemAlloc_v2@PLT
	movq	8(%rsp), %rsi
	movq	(%rsp), %rcx
	movabsq	$4890346246494415225, %rax
	movl	$16, %edx
	movq	%rax, 512(%rsi)
	addq	$512, %rsi
	movabsq	$4818288652467601408, %rax
	movq	%rax, 8(%rsi)
	movq	16(%rsp), %rax
	leaq	256(%rax), %rdi
	call	cuMemcpyHtoDAsync_v2@PLT
	movq	(%rsp), %rdi
	call	cuStreamSynchronize@PLT
	movq	8(%rsp), %rdi
	xorl	%eax, %eax
	movq	(%rsp), %rcx
	movq	16(%rsp), %rsi
	movl	$4096, %edx
	movq	%rax, 256(%rdi)
	movq	%rax, 264(%rdi)
	call	cuMemcpyDtoHAsync_v2@PLT
	movq	(%rsp), %rdi
	call	cuStreamSynchronize@PLT
	movq	8(%rsp), %rax
	movq	%rbx, %rdi
	cvtss2sd	256(%rax), %xmm0
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	%rax, %rdi
	call	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@PLT
	movq	8(%rsp), %rax
	movq	%rbx, %rdi
	cvtss2sd	260(%rax), %xmm0
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	%rax, %rdi
	call	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@PLT
	movq	8(%rsp), %rax
	movq	%rbx, %rdi
	cvtss2sd	264(%rax), %xmm0
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	%rax, %rdi
	call	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@PLT
	movq	8(%rsp), %rax
	movq	%rbx, %rdi
	cvtss2sd	268(%rax), %xmm0
	call	_ZNSo9_M_insertIdEERSoT_@PLT
	movq	%rax, %rdi
	call	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@PLT
	movq	8(%rsp), %rdi
	call	cuMemFreeHost@PLT
	movq	16(%rsp), %rdi
	call	cuMemFree_v2@PLT
	movq	(%rsp), %rdi
	call	cuStreamDestroy_v2@PLT
	movq	24(%rsp), %rax
	subq	%fs:40, %rax
	je	.L3
	call	__stack_chk_fail@PLT
.L3:
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	xorl	%eax, %eax
	popq	%rbx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4279:
	.size	main, .-main
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB4305:
	.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
.LFE4305:
	.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
	.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:

// What if we are given a device pointer that is offset from any of the device pointers we provided to the client?
//
// This file is a test-case for this.  Then we can look at handling that...

#include <iostream>
#include <memory>
#include <cassert>

using namespace std;

#include <hip/hip_runtime.h>

// __global__ void getValue(float *data, int idx, float value) {
//     if(threadIdx.x == 0) {
//         data[idx] += value;
//     }
// }

int main(int argc, char *argv[]) {
    int N = 1024;

    hipStream_t stream;
    hipStreamCreateWithFlags(&stream, 0);

    float *hostFloats;
    hipHostAlloc((void **)&hostFloats, N * sizeof(float), hipHostMallocPortable);

    hipDeviceptr_t deviceFloats;
    hipMalloc(&deviceFloats, N * sizeof(float));

    hostFloats[128] = 123.456f;
    hostFloats[129] = 444.0f;
    hostFloats[130] = 321.0f;
    hostFloats[131] = 111.0f;

    // now we will copy 16 bytes, starting at location 128...
    hipMemcpyHtoDAsync(
        (hipDeviceptr_t)(((float *)deviceFloats) + 64),
        hostFloats + 128,
        4 * sizeof(float),
        stream
    );
    hipStreamSynchronize(stream);
    // now copy back entire buffer
    hostFloats[64] = 0.0f;
    hostFloats[65] = 0.0f;
    hostFloats[66] = 0.0f;
    hostFloats[67] = 0.0f;
    hipMemcpyDtoHAsync(hostFloats, deviceFloats, N * sizeof(float), stream);
    hipStreamSynchronize(stream);

    // and check the values...
    cout << hostFloats[64] << endl;
    cout << hostFloats[65] << endl;
    cout << hostFloats[66] << endl;
    cout << hostFloats[67] << endl;

    assert(hostFloats[64] == 123.456f);
    assert(hostFloats[65] == 444.0f);
    assert(hostFloats[66] == 321);
    assert(hostFloats[67] == 111);

    hipHostFree(hostFloats);
    hipFree(deviceFloats);
    hipStreamDestroy(stream);

    return 0;
}

	.text
	.file	"testpartialcopy.hip"
                                        # Start of file scope inline assembly
	.globl	_ZSt21ios_base_library_initv

                                        # End of file scope inline assembly
	.section	.rodata.cst16,"aM",@progbits,16
	.p2align	4, 0x0                          # -- Begin function main
.LCPI0_0:
	.long	0x42f6e979                      # float 123.456001
	.long	0x43de0000                      # float 444
	.long	0x43a08000                      # float 321
	.long	0x42de0000                      # float 111
	.text
	.globl	main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r12
	.cfi_def_cfa_offset 32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	subq	$24, %rsp
	.cfi_def_cfa_offset 64
	.cfi_offset %rbx, -40
	.cfi_offset %r12, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	16(%rsp), %rbx
	movq	%rbx, %rdi
	xorl	%esi, %esi
	callq	hipStreamCreateWithFlags
	leaq	8(%rsp), %r14
	movl	$4096, %esi                     # imm = 0x1000
	movq	%r14, %rdi
	movl	$1, %edx
	callq	hipHostAlloc
	movq	%rsp, %r15
	movl	$4096, %esi                     # imm = 0x1000
	movq	%r15, %rdi
	callq	hipMalloc
	movq	(%r14), %rsi
	movaps	.LCPI0_0(%rip), %xmm0           # xmm0 = [1.23456001E+2,4.44E+2,3.21E+2,1.11E+2]
	movups	%xmm0, 512(%rsi)
	addq	$512, %rsi                      # imm = 0x200
	movl	$256, %edi                      # imm = 0x100
	addq	(%r15), %rdi
	movq	(%rbx), %rcx
	movl	$16, %edx
	callq	hipMemcpyHtoDAsync
	movq	(%rbx), %rdi
	callq	hipStreamSynchronize
	movq	(%r14), %rdi
	xorps	%xmm0, %xmm0
	movups	%xmm0, 256(%rdi)
	movq	(%r15), %rsi
	movq	(%rbx), %rcx
	movl	$4096, %edx                     # imm = 0x1000
	callq	hipMemcpyDtoHAsync
	movq	(%rbx), %rdi
	callq	hipStreamSynchronize
	movq	(%r14), %rax
	xorps	%xmm0, %xmm0
	cvtss2sd	256(%rax), %xmm0
	movl	$_ZSt4cout, %edi
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %r12
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%r12, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%r12, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movq	(%r14), %rax
	xorps	%xmm0, %xmm0
	cvtss2sd	260(%rax), %xmm0
	movl	$_ZSt4cout, %edi
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %r12
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%r12, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%r12, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movq	(%r14), %rax
	xorps	%xmm0, %xmm0
	cvtss2sd	264(%rax), %xmm0
	movl	$_ZSt4cout, %edi
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %r12
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%r12, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%r12, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movq	(%r14), %rax
	xorps	%xmm0, %xmm0
	cvtss2sd	268(%rax), %xmm0
	movl	$_ZSt4cout, %edi
	callq	_ZNSo9_M_insertIdEERSoT_
	movq	%rax, %r12
	movq	(%rax), %rax
	movq	-24(%rax), %rdi
	addq	%r12, %rdi
	movl	$10, %esi
	callq	_ZNKSt9basic_iosIcSt11char_traitsIcEE5widenEc
	movsbl	%al, %esi
	movq	%r12, %rdi
	callq	_ZNSo3putEc
	movq	%rax, %rdi
	callq	_ZNSo5flushEv
	movq	(%r14), %rdi
	callq	hipHostFree
	movq	(%r15), %rdi
	callq	hipFree
	movq	(%rbx), %rdi
	callq	hipStreamDestroy
	xorl	%eax, %eax
	addq	$24, %rsp
	.cfi_def_cfa_offset 40
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%r12
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end0:
	.size	main, .Lfunc_end0-main
	.cfi_endproc
                                        # -- End function
	.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 _ZSt4cout
	.addrsig_sym __hip_cuid_

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

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

	.end_amdgpu_metadata
