Outline

• Memory model and hierarchy
• Memory management strategies
• Page-locked memory
• Asynchronous memory allocation

Preface: Virtual Memory addressing

Memory model and hierarchy

Memory management strategies

Memory management can be Explicit or Implicit.

• HIP API documentation on memory

Memory management strategies

int main() {
 int *A, *d_A;
 A = (int *) malloc(N*sizeof(int));
 hipMalloc((void**)&d_A, N*sizeof(int));
 ...
 /* Copy data to GPU and launch kernel */
 hipMemcpy(d_A, A, N*sizeof(int), hipMemcpyHostToDevice);
 kernel<<<...>>>(d_A);
 hipMemcpy(A, d_A, N*sizeof(int), hipMemcpyDeviceToHost);
 hipFree(d_A);
 // result is in A
 free(A);
}

int main() {
 int *A;
 hipMallocManaged((void**)&A, N*sizeof(int));
 ...
 /* Launch GPU kernel */
 kernel<<<...>>>(A);
 hipStreamSynchronize(0);
 // result is in A
 hipFree(A);
}

Unified Memory pros & cons

Unified memory: Prefetching

• Unified memory automatically migrates memory pages between CPU and GPU

• Without prefetching:

  • Memory pages migrate on-demand
  • First GPU access may trigger page faults

• Programmer can proactively move pages to the GPU before execution

  hipError_t hipMemPrefetchAsync(void *dev_ptr, size_t size, int device, hipStream_t stream);

Explicit memory API calls

• Allocate device memory

  hipError_t hipMalloc(void **devPtr, size_t size)

• Copy data

  hipError_t hipMemcpy(void *dst, const void *src, size_t count, enum hipMemcpyKind kind)

  Where kind:

  • hipMemcpyDefault, or
    hipMemcpyDeviceToHost, hipMemcpyHostToDevice,
    hipMemcpyHostToHost, hipMemcpyDeviceToDevice

• Deallocate device memory

  hipError_t hipFree(void *devPtr)

Unified memory API calls

Also known as Managed memory

• Allocate Unified Memory

  hipError_t hipMallocManaged(void **devPtr, size_t size)

• Deallocate unified memory (same as explicitly managed memory)

  hipError_t hipFree(void *devPtr)

• Prefetch (asynchronously):

  hipError_t hipMemPrefetchAsync( void *dev_ptr, size_t size, int device, hipStream_t stream)

• Advise about memory access (more in HIP API Documentation)

  hipError_t hipMemAdvise(void *dev_ptr, size_t size, hipMemoryAdvise advise, int device)

Summary about memory management

• Memory management can be Explicit or Implicit
• Unified memory handles memory management between host and device “automatically”
• Any questions about explicit vs. implicit memory management?

Pageable vs. page-locked memory?

Page-locked (or pinned) memory

Asynchronous memcopies

• Normal hipMemcpy() calls are blocking (ie, synchronizing)
  • The execution of host code is blocked until copying is finished
• To overlap copying and program execution, use asynchronous functions
  • Such functions have Async suffix, eg, hipMemcpyAsync()
• User has to synchronize the program execution
• Concurrency with memory copy and computation requires page-locked host allocations

Async memory copy with regular vs page-locked memory

Explicit memory API calls

Page-locked host memory

• Allocate/free page-locked host memory

    hipHostMalloc(void **ptr, size_t size);
    hipHostFree(void *ptr);

• Memory copy functions are the same as with normally allocated memory

Asynchronous allocation: The stream-ordered memory allocator and memory pools

• Benefit of asynchronous memory management: allocate/free memory from/to a pool

Memory pools - Example

for (int i = 0; i < 100; i++) {
  // Allocate memory here (slow)
  hipMalloc(&ptr, size); 
  // Run GPU kernel
  kernel<<<..., stream>>>(ptr);
  // Deallocate memory here
  hipFree(ptr); 
}
// Synchronize the default stream (no influence to memory allocations)
hipStreamSynchronize(0); 

for (int i = 0; i < 100; i++) {
  // Obtain unused memory from the current memory pool, 
  // more memory is allocated for the pool if needed
  hipMallocAsync(&ptr, size, stream); 
  // Run GPU kernel
  kernel<<<..., stream>>>(ptr);
  // Return memory to the current memory pool
  hipFreeAsync(ptr, stream); 
}
// Synchronize 
hipStreamSynchronize(stream); 

Summary

• Host and device have separate physical memories
  • The data copies between CPU and GPU should be minimized
• Explicit or implicit memory management
  • Unified Memory: improve productivity and cleaner implementation
• Page-locked host allocation: DMA and kernel access to host memory
• Asynchronous allocation and deallocation: memory pools