unified.hip File Reference

#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <stdint.h>
#include <hip/hip_runtime.h>
#include <device/hip/unified.h>

Go to the source code of this file.

Functions
__global__ void	hip_unified_memset_kernel (unsigned char *a, unsigned char v, size_t n)
__global__ void	hip_unified_copy8_kernel (unsigned long long *__restrict__ a, const unsigned long long *__restrict__ b, size_t n)
__global__ void	hip_unified_copy1_kernel (unsigned char *__restrict__ a, const unsigned char *__restrict__ b, size_t n)
int	hip_zerocopy (void)
hipError_t	hip_map (void **ptr_d, void *ptr_h, size_t s)
static int	hip_is_device_alloc (void *ptr_d)
hipError_t	hip_map_free (void *ptr_d)
hipError_t	hip_map_memset (void *ptr_d, int value, size_t s, void *stream)
hipError_t	hip_map_memcpy (void *dst, void *src, size_t s, int kind, void *stream)

Function Documentation

hip_is_device_alloc()

static int hip_is_device_alloc ( void *  ptr_d )

static

Whether ptr_d is a device allocation (hipMalloc); zero-copy mappings are unregistered host pointers.

Definition at line 205 of file unified.hip.

Here is the call graph for this function:

Here is the caller graph for this function:

hip_map()

hipError_t hip_map	(	void **	ptr_d,
		void *	ptr_h,
		size_t	s
	)

Map s bytes of host memory at ptr_h to the device.

On unified memory (see hip_zerocopy) the device pointer aliases the host allocation; otherwise a separate replicated buffer is allocated with hipMalloc. Aliased ranges are advised to coarse-grained coherence, since mapped arrays are never accessed concurrently from host and device and fine-grained system memory can be slower in kernels.

Definition at line 165 of file unified.hip.

Here is the call graph for this function:

hip_map_free()

hipError_t hip_map_free

(

void *

ptr_d

)

Free a device pointer obtained from hip_map (or hipMalloc).

Pointers aliasing host memory are left untouched; their allocation is owned by the host side. hipFree implicitly synchronizes the device, and the host typically deallocates right after an unmap, so preserve that barrier for aliased mappings such that the allocation cannot be released under in-flight device work.

Definition at line 228 of file unified.hip.

Here is the call graph for this function:

hip_map_memcpy()

hipError_t hip_map_memcpy	(	void *	dst,
		void *	src,
		size_t	s,
		int	kind,
		void *	stream
	)

Memcpy between device pointers obtained from hip_map (or hipMalloc) and/or host pointers.

Under zero-copy either side may alias pageable host memory, which hipMemcpy treats as a staged pageable copy (slow, serializing); copy with a kernel instead, which runs at full memory bandwidth and stays stream-ordered (on a unified memory APU any pointer is dereferenceable from the device).

Definition at line 269 of file unified.hip.

Here is the call graph for this function:

hip_map_memset()

hipError_t hip_map_memset	(	void *	ptr_d,
		int	value,
		size_t	s,
		void *	stream
	)

Memset on a device pointer obtained from hip_map (or hipMalloc).

hipMemset only accepts device allocations, so pointers aliasing host memory are set with a kernel instead, which also keeps the operation stream-ordered.

Definition at line 243 of file unified.hip.

Here is the call graph for this function:

hip_unified_copy1_kernel()

__global__ void hip_unified_copy1_kernel	(	unsigned char *__restrict__	a,
		const unsigned char *__restrict__	b,
		size_t	n
	)

Definition at line 83 of file unified.hip.

Here is the call graph for this function:

Here is the caller graph for this function:

hip_unified_copy8_kernel()

__global__ void hip_unified_copy8_kernel	(	unsigned long long *__restrict__	a,
		const unsigned long long *__restrict__	b,
		size_t	n
	)

Grid-stride copy kernels for zero-copy mappings (word and byte variants).

Definition at line 73 of file unified.hip.

Here is the call graph for this function:

Here is the caller graph for this function:

hip_unified_memset_kernel()

__global__ void hip_unified_memset_kernel	(	unsigned char *	a,
		unsigned char	v,
		size_t	n
	)

Zero-copy mapping of host arrays for unified memory architectures.

On APUs with a single physical memory (e.g. MI300A) and XNACK enabled (HSA_XNACK=1), system allocations are directly and coherently accessible from device kernels. Mapped host arrays then alias their host allocation instead of being replicated in a separate hipMalloc buffer, and host-device copies degenerate to no-ops. Grid-stride byte memset kernel for zero-copy mappings.

Definition at line 60 of file unified.hip.

Here is the call graph for this function:

Here is the caller graph for this function:

hip_zerocopy()

int hip_zerocopy

(

void

)

Whether mapped arrays may alias host allocations (zero-copy).

Zero-copy is OFF by default and must be opted into with NEKO_HIP_ZEROCOPY=1: it is correct on a unified-memory APU but currently slower than replicated buffers on MI300A (the gather- scatter exchange in particular), so replication is the default until that overhead is understood. It also halves the memory footprint of mapped data, which is its reason to exist as an opt-in capacity mode.

When opted in, zero-copy is still only possible on an APU with unified physical memory and XNACK enabled, such that system allocations are directly and coherently accessible from the device:

• The integrated attribute identifies an APU (e.g. MI300A, as opposed to the discrete MI300X, which shares the same gfx942 architecture). Discrete GPUs (MI250X, MI300X) always use replicated buffers; with XNACK they also report pageable memory access, but there it means page migration, which would thrash host arrays back and forth.
• Pageable memory access reflects the effective XNACK state of the process (HSA_XNACK=1, kernel support and xnack-capable code objects); on an APU with XNACK off it reports 0 and we fall back to replicated buffers.

Definition at line 122 of file unified.hip.

Here is the call graph for this function:

Here is the caller graph for this function: