[Issue]: MI300A : hipMalloc falls in wrong GPU HBM under memory pressure

[cboillot]

Problem Description

This issue arises when using multiple MI300A APUs in various scenarios, such as during high GPU memory usage or after loading AI model tensors. We observe a huge performance drop (by a factor of 2 to 10), primarily because some of the hipMalloc allocations on one device are being partially fulfilled by the HBM of another device. We expectthat all allocations on a GPU device should be served exclusively by the memory of that selected device, or the allocation must fail.

Operating System

Red Hat Enterprise Linux 9.4 (Plow)

CPU

AMD Instinct MI300A Accelerator

GPU

4 * AMD Instinct MI300A Accelerator

ROCm Version

ROCm 6.4.0

ROCm Component

No response

Steps to Reproduce

To reproduce we used a python env with torch 2.7.0 and numpy (reproducer.py is below), but you can reproduce using hipMalloc by monitoring the memory usage repartition in /sys/devices/system/node/node*/meminfo.
python reproducer.py --first_gpu_alloc_ratio 0.95 --next_gpus_relative_alloc_ratio 0.90

Output :

Number of available GPUs: 4
-------------------------------
Test 1: This pass will execute the following steps:
1. Show the current memory usage on each NUMA nodes.
2. Allocate tensors on GPU 1 to fill its memory capacity.
3. Check the actual location of the allocated memory by checking in which NUMA node's memory usage has increased.
4. Evaluate TFLOPs by selecting 3 random tensors (A, B, C) on the GPU and compute C += A . B (dot product).
5. Repeat the same process on each GPU without releasing the memory from the previous GPU.
--
Free memory layout at startup :
Numa node 0 free memory : 121 GiB, (of which pagecache memory : 1 GiB)
Numa node 1 free memory : 124 GiB, (of which pagecache memory : 0 GiB)
Numa node 2 free memory : 124 GiB, (of which pagecache memory : 0 GiB)
Numa node 3 free memory : 124 GiB, (of which pagecache memory : 0 GiB)
--
GPU 1 : alloc 95.0% - 121.6 GiB of the GPU memory in 7782 tensors
 Allocation tooks 13 seconds and comes at 93% (10.09 GiB still free of which 0.45 GiB of pagecache) from the right numa and from :
  NUMA node 2: 7.72 GiB
 Bench (any bank) :: 23.666840 TFLOPs (during 1.451455 sec)
 Bench (right bank) :: 28.834836 TFLOPs (during 1.191314 sec)
 Bench (wrong bank) :: 5.774188 TFLOPs (during 5.949122 sec)
--
GPU 2 : alloc 85.5% - 109.44 GiB of the GPU memory in 7004 tensors
 Allocation tooks 13 seconds and comes mainly from the right memory bank
 Bench (any bank) :: 54.405166 TFLOPs (during 0.631399 sec)
 Bench (right bank) :: 54.367535 TFLOPs (during 0.631836 sec)
--
GPU 3 : alloc 76.95% - 98.5 GiB of the GPU memory in 6303 tensors
 Allocation tooks 12 seconds and comes mainly from the right memory bank
 Bench (any bank) :: 54.410138 TFLOPs (during 0.631341 sec)
 Bench (right bank) :: 54.393623 TFLOPs (during 0.631533 sec)
--
GPU 0 : alloc 69.25% - 88.65 GiB of the GPU memory in 5673 tensors
 Allocation tooks 6 seconds and comes mainly from the right memory bank
 Bench (any bank) :: 54.396970 TFLOPs (during 0.631494 sec)
 Bench (right bank) :: 54.373855 TFLOPs (during 0.631762 sec)


-------------------------------
Test 2: This pass will execute the following steps:
1. Fill the pagecache with a dummy 100 GiB file to demonstrate the effect of pagecache on memory allocation.
2. Retry the allocation and performance evaluation after partially filling the pagecache.
107374182400 bytes (107 GB, 100 GiB) copied, 102 s, 1.1 GB/s
100+0 records in
100+0 records out
107374182400 bytes (107 GB, 100 GiB) copied, 101.898 s, 1.1 GB/s
Successfully created 100 GiB file filled with zeros: ./zero_file.bin
Successfully synchronized file system buffers.
--
Free memory layout at startup :
Numa node 0 free memory : 111 GiB, (of which pagecache memory : 101 GiB)
Numa node 1 free memory : 124 GiB, (of which pagecache memory : 0 GiB)
Numa node 2 free memory : 124 GiB, (of which pagecache memory : 0 GiB)
Numa node 3 free memory : 124 GiB, (of which pagecache memory : 0 GiB)
--
GPU 1 : alloc 95.0% - 121.6 GiB of the GPU memory in 7782 tensors
 Allocation tooks 13 seconds and comes at 93% (9.93 GiB still free of which 0.45 GiB of pagecache) from the right numa and from :
  NUMA node 2: 7.55 GiB
 Bench (any bank) :: 24.073996 TFLOPs (during 1.426907 sec)
 Bench (right bank) :: 28.760786 TFLOPs (during 1.194381 sec)
 Bench (wrong bank) :: 5.542506 TFLOPs (during 6.197801 sec)
--
GPU 2 : alloc 85.5% - 109.44 GiB of the GPU memory in 7004 tensors
 Allocation tooks 13 seconds and comes mainly from the right memory bank
 Bench (any bank) :: 54.412645 TFLOPs (during 0.631312 sec)
 Bench (right bank) :: 54.356111 TFLOPs (during 0.631968 sec)
--
GPU 3 : alloc 76.95% - 98.5 GiB of the GPU memory in 6303 tensors
 Allocation tooks 11 seconds and comes mainly from the right memory bank
 Bench (any bank) :: 54.441244 TFLOPs (during 0.630980 sec)
 Bench (right bank) :: 54.393767 TFLOPs (during 0.631531 sec)
--
GPU 0 : alloc 69.25% - 88.65 GiB of the GPU memory in 5673 tensors
 Allocation tooks 67 seconds and comes at 75% (49.43 GiB still free of which 0.41 GiB of pagecache) from the right numa and from :
  NUMA node 3: 19.7 GiB
 Bench (any bank) :: 22.395341 TFLOPs (during 1.533861 sec)
 Bench (right bank) :: 18.161142 TFLOPs (during 1.891475 sec)
 Bench (wrong bank) :: 54.232145 TFLOPs (during 0.633413 sec)


Successfully removed: ./zero_file.bin

reproducer.py :

import os
import torch
import subprocess
import time
import random
import argparse

def main(args):

    # Get the number of available GPUs
    num_gpus = torch.cuda.device_count()
    print(f"Number of available GPUs: {num_gpus}")

    # Execute 2 pass of :
    # --> Show the memory 
    #  --> allocate tensors on GPU 1 to fill it's memory
    #  --> Check the real location of the allocated memory (which NUMA node memory increased ?)
    #  --> Evaluate TFLOPs by selecting 3 random tensors (A, B, C) in the GPU and compute C += A . B (dot product)
    #  ... same thing on each of the GPU (without releasing the memory of each previous GPU)
    #
    # the second pass differ only by the fact that we start by partially filling pagecache 
    # with evictable pages to show the effect of the pagecache on 
    for l in range(2) :
        alloc_ratio = args.first_gpu_alloc_ratio
        torch.cuda.empty_cache()
        print("-------------------------------")
        if l == 0 :
            print(f"Test 1: This pass will execute the following steps:")
            print("1. Show the current memory usage on each NUMA nodes.")
            print(f"2. Allocate tensors on GPU {args.first_gpu_index % num_gpus} to fill its memory capacity.")
            print("3. Check the actual location of the allocated memory by checking in which NUMA node's memory usage has increased.")
            print("4. Evaluate TFLOPs by selecting 3 random tensors (A, B, C) on the GPU and compute C += A . B (dot product).")
            print("5. Repeat the same process on each GPU without releasing the memory from the previous GPU.")
        else :
            if not args.do_pagecache_test :
                break
            print(f"Test 2: This pass will execute the following steps:")
            print(f"1. Fill the pagecache with a dummy {args.pagecache_fill_size} GiB file to demonstrate the effect of pagecache on memory allocation.")
            print("2. Retry the allocation and performance evaluation after partially filling the pagecache.")
            create_zero_file('./zero_file.bin', args.pagecache_fill_size)
        print("--")

        print("Free memory layout at startup :")
        memfree_start, pagecache_start = get_memfree()
        for node_id, memfree_bytes in memfree_start.items() :
            print(f"Numa node {node_id} free memory : {round(memfree_bytes/1024/1024)} GiB, (of which pagecache memory : {round(pagecache_start[node_id]/1024/1024)} GiB)")
        tensors = [None] * num_gpus


        for i in range(num_gpus) :
            print("--")
            device_idx = (i + args.first_gpu_index) % num_gpus
            total_memory = torch.cuda.get_device_properties(device_idx).total_memory

            # Get initial numa node memory info
            memfree_before, pagecache_before = get_memfree()
            
            # Allocate a bunch of tensors (nb_blocks) that fill alloc_ratio of the GPU
            nb_blocks = round((total_memory * alloc_ratio) // (args.matrix_size * args.matrix_size * 4))

            print(f"GPU {device_idx} : alloc {round(alloc_ratio * 100, 2)}% - {round(total_memory*alloc_ratio / 1024 ** 3, 2)} GiB of the GPU memory in {nb_blocks} tensors")
            start_time = time.time()
            tensors[device_idx] = [torch.zeros(args.matrix_size, args.matrix_size, dtype=torch.float32, device=f"cuda:{device_idx}") for j in range(nb_blocks)]
            torch.cuda.synchronize(device_idx)
            duration = time.time() - start_time
            
            # Get numa node memory info after allocation
            memfree_after, pagecache_after = get_memfree()

            # Calculate the difference in memory 
            mem_diff = {}
            total_diff = 0
            secondary_pools = ""
            for node in memfree_before.keys():
                diff = memfree_after[node] - memfree_before[node]
                mem_diff[node] = diff
                total_diff += diff
                if node != device_idx and diff < - 1024 * 1024:
                    # Print additional nodes where the allocation difference is greater than 1024 MiB -- which is enough to skip the python memory usage increase
                    secondary_pools += f"\n  NUMA node {node}: {-round(diff/1024/1024, 2)} GiB"

            ratio_in_right_numa = mem_diff[device_idx] / total_diff

            if len(secondary_pools) == 0 :
                print(f" Allocation tooks {round(duration)} seconds and comes mainly from the right memory bank")
            else :
                print(f" Allocation tooks {round(duration)} seconds and comes at {round(ratio_in_right_numa*100)}% ({round(memfree_after[device_idx] / 1024 / 1024, 2)} GiB still free of which {round(pagecache_after[device_idx]/1024/1024, 2)} GiB of pagecache) from the right numa and from :{secondary_pools}")

            # Perform matrix multiplication on each GPU
            if args.do_perf_test :
                execute_perf_test(args, tensors, device_idx, nb_blocks, ratio_in_right_numa)
            
            # reduce the % of the memory to show different behaviours
            alloc_ratio *= args.next_gpus_relative_alloc_ratio
        tensors = None
        print()
        print()

    remove_file('./zero_file.bin')

# Get the free and pagecache memory for each numa node from /sys/devices/system/node/node*/meminfo.
def get_memfree():
    memfree = {}
    page_cache = {}
    try:
        output = subprocess.check_output("grep -hi 'MemFree' /sys/devices/system/node/node*/meminfo", shell=True)
        for line in output.decode().strip().split('\n'):
            # assume format is : Node 0 MemFree:        28046496 kB 
            parts = line.split()
            node_id = int(parts[1])  
            memfree[node_id] = int(parts[3])

        output = subprocess.check_output("grep -hi 'FilePages' /sys/devices/system/node/node*/meminfo", shell=True)
        for line in output.decode().strip().split('\n'):
            # assume format is : Node 0 MemFree:        28046496 kB 
            parts = line.split()
            node_id = int(parts[1])  
            page_cache[node_id] = int(parts[3])
            memfree[node_id] += page_cache[node_id]

    except Exception as e:
        print(f"Error reading memory info: {e}")
    return memfree, page_cache

def pick_random_square_matrix(tensors, mnp, n_blocks) : 
    if n_blocks < 0 :
        idx = random.randint(n_blocks, -1) # pick a random index for tensor
    else :
        idx = random.randint(0, n_blocks - 1) # pick a random index for tensor
    return tensors[idx].view(mnp, mnp) # use the tensor as a square matrix


def random_mm_sum(tensors, mnp, n_blocks, n_mm) :
    for l in range(n_mm):
        A = pick_random_square_matrix(tensors, mnp, n_blocks)
        B = pick_random_square_matrix(tensors, mnp, n_blocks)
        C = pick_random_square_matrix(tensors, mnp, n_blocks)
        C += torch.mm(A,B)
    flops = n_mm * (mnp * mnp * (2*mnp - 1)) 
    return C, flops

def execute_perf_test(args, tensors, device_idx, nb_blocks, ratio_in_right_numa) :
    results = []
    def bench(test_name, n_blocks, n_mm, from_end = False) :
        torch.cuda.synchronize(device=f'cuda:{device_idx}')
        start_time = time.time()
        C, total_flops = random_mm_sum(tensors[device_idx], args.matrix_size, n_blocks, n_mm)
        torch.cuda.synchronize(device=f'cuda:{device_idx}')
        end_time = time.time()

        elapsed_time = end_time - start_time  # Time in seconds
        tflops = total_flops / (elapsed_time * 1e12)  # Convert to TFLOPs
        if test_name is not None :
            print(f"{test_name}: {tflops:.6f} TFLOPs (during {elapsed_time:.6f} sec)")
        return C

    # warmup bench (discarded)
    results.append(bench(None, nb_blocks, args.perf_nb_warmup_loop))

    # any tensor is candidate
    results.append(bench(" Bench (any bank) :", nb_blocks, args.perf_nb_loop))
    
    # assume that the allocation started to use the right memory bank and then fallback to the wrong one
    nb_block_in_right_numa = round(ratio_in_right_numa * nb_blocks) - 1
    if nb_block_in_right_numa > 100 :
        results.append(bench(" Bench (right bank) :", nb_block_in_right_numa, args.perf_nb_loop))

    nb_block_in_wrong_numa = round((1 - ratio_in_right_numa) * nb_blocks) - 1
    if nb_block_in_wrong_numa > 100 :
        results.append(bench(" Bench (wrong bank) :", -nb_block_in_wrong_numa, args.perf_nb_loop))


# Write a file full of zeros that will fill the pagecache
# you can replace it by any non O_DIRECT file read (for instance cat filename > /dev/null)
# we ensure the file is properly sync at the end to be sure pages can be evicted from pagecache easily
def create_zero_file(filename, size_gb):
    # Execute the dd command
    process = subprocess.run(f'dd if=/dev/zero of={filename} bs=1G count={size_gb} status=progress', shell=True)

    # Check if the command was successful
    if process.returncode == 0:
        print(f"Successfully created {size_gb} GiB file filled with zeros: {filename}")
    else:
        print(f"Error occurred while creating the file: {process.returncode}")
    
    # Ensure there are no dirty pages
    sync_process = subprocess.run('sync', shell=True)
    if sync_process.returncode == 0:
        print("Successfully synchronized file system buffers.")
    else:
        print(f"Error occurred while synchronizing: {sync_process.returncode}")


# At the end we want to empty the pagecache to restore the original state
# it is done by destroying the file created by create_zero_file
def remove_file(filename):
    try:
        os.remove(filename)
        print(f"Successfully removed: {filename}")
    except FileNotFoundError:
        print(f"File not found: {filename}")
    except PermissionError:
        print(f"Permission denied: {filename}")
    except Exception as e:
        print(f"Error occurred while removing the file: {e}")


if __name__ == "__main__":
    parser = argparse.ArgumentParser(description="GPU Memory Allocation Parameters")

    parser.add_argument('--first_gpu_alloc_ratio', type=float, required=True,
                        help='Percent of the GPU capacity to allocate on the first GPU')
    parser.add_argument('--next_gpus_relative_alloc_ratio', type=float, required=True,
                        help='Each GPU allocates memory of the previous one multiplied by this ratio')
    parser.add_argument('--first_gpu_index', type=int, default=1,
                        help='Index to start with another GPU than the 0')
    parser.add_argument('--matrix_size', type=int, default=2048,
                        help='Size of the tensors to allocate (matrix_size x matrix_size)')

    parser.add_argument('--do_perf_test', type=bool, default=True,
                        help='Flag to enable performance test')
    parser.add_argument('--perf_nb_loop', type=int, default=2000,
                        help='Number of dot products with 2 matrices of shape (matrix_size, matrix_size)')
    parser.add_argument('--perf_nb_warmup_loop', type=int, default=round(2000 * 0.2),
                        help='Number of warmup loops for performance test')

    parser.add_argument('--do_pagecache_test', type=bool, default=True,
                        help='Flag to enable pagecache test')
    parser.add_argument('--pagecache_fill_size', type=int, default=100,
                        help='Size in GiB to fill in the pagecache for the test')

    args = parser.parse_args()
    main(args)

### (Optional for Linux users) Output of /opt/rocm/bin/rocminfo --support

ROCk module version 6.7.0 is loaded
=====================    
HSA System Attributes    
=====================    
Runtime Version:         1.15
Runtime Ext Version:     1.7
System Timestamp Freq.:  1000.000000MHz
Sig. Max Wait Duration:  18446744073709551615 (0xFFFFFFFFFFFFFFFF) (timestamp count)
Machine Model:           LARGE                              
System Endianness:       LITTLE                             
Mwaitx:                  DISABLED
XNACK enabled:           YES
DMAbuf Support:          YES
VMM Support:             YES

==========               
HSA Agents               
==========               
*******                  
Agent 1                  
*******                  
  Name:                    AMD Instinct MI300A Accelerator    
  Uuid:                    CPU-XX                             
  Marketing Name:          AMD Instinct MI300A Accelerator    
  Vendor Name:             CPU                                
  Feature:                 None specified                     
  Profile:                 FULL_PROFILE                       
  Float Round Mode:        NEAR                               
  Max Queue Number:        0(0x0)                             
  Queue Min Size:          0(0x0)                             
  Queue Max Size:          0(0x0)                             
  Queue Type:              MULTI                              
  Node:                    0                                  
  Device Type:             CPU                                
  Cache Info:              
    L1:                      32768(0x8000) KB                   
  Chip ID:                 0(0x0)                             
  ASIC Revision:           0(0x0)                             
  Cacheline Size:          64(0x40)                           
  Max Clock Freq. (MHz):   3700                               
  BDFID:                   0                                  
  Internal Node ID:        0                                  
  Compute Unit:            48                                 
  SIMDs per CU:            0                                  
  Shader Engines:          0                                  
  Shader Arrs. per Eng.:   0                                  
  WatchPts on Addr. Ranges:1                                  
  Memory Properties:       
  Features:                None
  Pool Info:               
    Pool 1                   
      Segment:                 GLOBAL; FLAGS: FINE GRAINED        
      Size:                    130847776(0x7cc9420) KB            
      Allocatable:             TRUE                               
      Alloc Granule:           4KB                                
      Alloc Recommended Granule:4KB                                
      Alloc Alignment:         4KB                                
      Accessible by all:       TRUE                               
    Pool 2                   
      Segment:                 GLOBAL; FLAGS: EXTENDED FINE GRAINED
      Size:                    130847776(0x7cc9420) KB            
      Allocatable:             TRUE                               
      Alloc Granule:           4KB                                
      Alloc Recommended Granule:4KB                                
      Alloc Alignment:         4KB                                
      Accessible by all:       TRUE                               
    Pool 3                   
      Segment:                 GLOBAL; FLAGS: KERNARG, FINE GRAINED
      Size:                    130847776(0x7cc9420) KB            
      Allocatable:             TRUE                               
      Alloc Granule:           4KB                                
      Alloc Recommended Granule:4KB                                
      Alloc Alignment:         4KB                                
      Accessible by all:       TRUE                               
    Pool 4                   
      Segment:                 GLOBAL; FLAGS: COARSE GRAINED      
      Size:                    130847776(0x7cc9420) KB            
      Allocatable:             TRUE                               
      Alloc Granule:           4KB                                
      Alloc Recommended Granule:4KB                                
      Alloc Alignment:         4KB                                
      Accessible by all:       TRUE                               
  ISA Info:                
*******                  
Agent 2                  
*******                  
  Name:                    AMD Instinct MI300A Accelerator    
  Uuid:                    CPU-XX                             
  Marketing Name:          AMD Instinct MI300A Accelerator    
  Vendor Name:             CPU                                
  Feature:                 None specified                     
  Profile:                 FULL_PROFILE                       
  Float Round Mode:        NEAR                               
  Max Queue Number:        0(0x0)                             
  Queue Min Size:          0(0x0)                             
  Queue Max Size:          0(0x0)                             
  Queue Type:              MULTI                              
  Node:                    1                                  
  Device Type:             CPU                                
  Cache Info:              
    L1:                      32768(0x8000) KB                   
  Chip ID:                 0(0x0)                             
  ASIC Revision:           0(0x0)                             
  Cacheline Size:          64(0x40)                           
  Max Clock Freq. (MHz):   3700                               
  BDFID:                   0                                  
  Internal Node ID:        1                                  
  Compute Unit:            48                                 
  SIMDs per CU:            0                                  
  Shader Engines:          0                                  
  Shader Arrs. per Eng.:   0                                  
  WatchPts on Addr. Ranges:1                                  
  Memory Properties:       
  Features:                None
  Pool Info:               
    Pool 1                   
      Segment:                 GLOBAL; FLAGS: FINE GRAINED        
      Size:                    131809200(0x7db3fb0) KB            
      Allocatable:             TRUE                               
      Alloc Granule:           4KB                                
      Alloc Recommended Granule:4KB                                
      Alloc Alignment:         4KB                                
      Accessible by all:       TRUE                               
    Pool 2                   
      Segment:                 GLOBAL; FLAGS: EXTENDED FINE GRAINED
      Size:                    131809200(0x7db3fb0) KB            
      Allocatable:             TRUE                               
      Alloc Granule:           4KB                                
      Alloc Recommended Granule:4KB                                
      Alloc Alignment:         4KB                                
      Accessible by all:       TRUE                               
    Pool 3                   
      Segment:                 GLOBAL; FLAGS: KERNARG, FINE GRAINED
      Size:                    131809200(0x7db3fb0) KB            
      Allocatable:             TRUE                               
      Alloc Granule:           4KB                                
      Alloc Recommended Granule:4KB                                
      Alloc Alignment:         4KB                                
      Accessible by all:       TRUE                               
    Pool 4                   
      Segment:                 GLOBAL; FLAGS: COARSE GRAINED      
      Size:                    131809200(0x7db3fb0) KB            
      Allocatable:             TRUE                               
      Alloc Granule:           4KB                                
      Alloc Recommended Granule:4KB                                
      Alloc Alignment:         4KB                                
      Accessible by all:       TRUE                               
  ISA Info:                
*******                  
Agent 3                  
*******                  
  Name:                    AMD Instinct MI300A Accelerator    
  Uuid:                    CPU-XX                             
  Marketing Name:          AMD Instinct MI300A Accelerator    
  Vendor Name:             CPU                                
  Feature:                 None specified                     
  Profile:                 FULL_PROFILE                       
  Float Round Mode:        NEAR                               
  Max Queue Number:        0(0x0)                             
  Queue Min Size:          0(0x0)                             
  Queue Max Size:          0(0x0)                             
  Queue Type:              MULTI                              
  Node:                    2                                  
  Device Type:             CPU                                
  Cache Info:              
    L1:                      32768(0x8000) KB                   
  Chip ID:                 0(0x0)                             
  ASIC Revision:           0(0x0)                             
  Cacheline Size:          64(0x40)                           
  Max Clock Freq. (MHz):   3700                               
  BDFID:                   0                                  
  Internal Node ID:        2                                  
  Compute Unit:            48                                 
  SIMDs per CU:            0                                  
  Shader Engines:          0                                  
  Shader Arrs. per Eng.:   0                                  
  WatchPts on Addr. Ranges:1                                  
  Memory Properties:       
  Features:                None
  Pool Info:               
    Pool 1                   
      Segment:                 GLOBAL; FLAGS: FINE GRAINED        
      Size:                    131809208(0x7db3fb8) KB            
      Allocatable:             TRUE                               
      Alloc Granule:           4KB                                
      Alloc Recommended Granule:4KB                                
      Alloc Alignment:         4KB                                
      Accessible by all:       TRUE                               
    Pool 2                   
      Segment:                 GLOBAL; FLAGS: EXTENDED FINE GRAINED
      Size:                    131809208(0x7db3fb8) KB            
      Allocatable:             TRUE                               
      Alloc Granule:           4KB                                
      Alloc Recommended Granule:4KB                                
      Alloc Alignment:         4KB                                
      Accessible by all:       TRUE                               
    Pool 3                   
      Segment:                 GLOBAL; FLAGS: KERNARG, FINE GRAINED
      Size:                    131809208(0x7db3fb8) KB            
      Allocatable:             TRUE                               
      Alloc Granule:           4KB                                
      Alloc Recommended Granule:4KB                                
      Alloc Alignment:         4KB                                
      Accessible by all:       TRUE                               
    Pool 4                   
      Segment:                 GLOBAL; FLAGS: COARSE GRAINED      
      Size:                    131809208(0x7db3fb8) KB            
      Allocatable:             TRUE                               
      Alloc Granule:           4KB                                
      Alloc Recommended Granule:4KB                                
      Alloc Alignment:         4KB                                
      Accessible by all:       TRUE                               
  ISA Info:                
*******                  
Agent 4                  
*******                  
  Name:                    AMD Instinct MI300A Accelerator    
  Uuid:                    CPU-XX                             
  Marketing Name:          AMD Instinct MI300A Accelerator    
  Vendor Name:             CPU                                
  Feature:                 None specified                     
  Profile:                 FULL_PROFILE                       
  Float Round Mode:        NEAR                               
  Max Queue Number:        0(0x0)                             
  Queue Min Size:          0(0x0)                             
  Queue Max Size:          0(0x0)                             
  Queue Type:              MULTI                              
  Node:                    3                                  
  Device Type:             CPU                                
  Cache Info:              
    L1:                      32768(0x8000) KB                   
  Chip ID:                 0(0x0)                             
  ASIC Revision:           0(0x0)                             
  Cacheline Size:          64(0x40)                           
  Max Clock Freq. (MHz):   3700                               
  BDFID:                   0                                  
  Internal Node ID:        3                                  
  Compute Unit:            48                                 
  SIMDs per CU:            0                                  
  Shader Engines:          0                                  
  Shader Arrs. per Eng.:   0                                  
  WatchPts on Addr. Ranges:1                                  
  Memory Properties:       
  Features:                None
  Pool Info:               
    Pool 1                   
      Segment:                 GLOBAL; FLAGS: FINE GRAINED        
      Size:                    131795632(0x7db0ab0) KB            
      Allocatable:             TRUE                               
      Alloc Granule:           4KB                                
      Alloc Recommended Granule:4KB                                
      Alloc Alignment:         4KB                                
      Accessible by all:       TRUE                               
    Pool 2                   
      Segment:                 GLOBAL; FLAGS: EXTENDED FINE GRAINED
      Size:                    131795632(0x7db0ab0) KB            
      Allocatable:             TRUE                               
      Alloc Granule:           4KB                                
      Alloc Recommended Granule:4KB                                
      Alloc Alignment:         4KB                                
      Accessible by all:       TRUE                               
    Pool 3                   
      Segment:                 GLOBAL; FLAGS: KERNARG, FINE GRAINED
      Size:                    131795632(0x7db0ab0) KB            
      Allocatable:             TRUE                               
      Alloc Granule:           4KB                                
      Alloc Recommended Granule:4KB                                
      Alloc Alignment:         4KB                                
      Accessible by all:       TRUE                               
    Pool 4                   
      Segment:                 GLOBAL; FLAGS: COARSE GRAINED      
      Size:                    131795632(0x7db0ab0) KB            
      Allocatable:             TRUE                               
      Alloc Granule:           4KB                                
      Alloc Recommended Granule:4KB                                
      Alloc Alignment:         4KB                                
      Accessible by all:       TRUE                               
  ISA Info:                
*******                  
Agent 5                  
*******                  
  Name:                    gfx942                             
  Uuid:                    GPU-6cdb410c2f219ac0               
  Marketing Name:          AMD Instinct MI300A                
  Vendor Name:             AMD                                
  Feature:                 KERNEL_DISPATCH                    
  Profile:                 BASE_PROFILE                       
  Float Round Mode:        NEAR                               
  Max Queue Number:        128(0x80)                          
  Queue Min Size:          64(0x40)                           
  Queue Max Size:          131072(0x20000)                    
  Queue Type:              MULTI                              
  Node:                    4                                  
  Device Type:             GPU                                
  Cache Info:              
    L1:                      32(0x20) KB                        
    L2:                      24576(0x6000) KB                   
    L3:                      262144(0x40000) KB                 
  Chip ID:                 29856(0x74a0)                      
  ASIC Revision:           1(0x1)                             
  Cacheline Size:          64(0x40)                           
  Max Clock Freq. (MHz):   2100                               
  BDFID:                   512                                
  Internal Node ID:        4                                  
  Compute Unit:            228                                
  SIMDs per CU:            4                                  
  Shader Engines:          24                                 
  Shader Arrs. per Eng.:   1                                  
  WatchPts on Addr. Ranges:4                                  
  Coherent Host Access:    TRUE                               
  Memory Properties:       APU
  Features:                KERNEL_DISPATCH 
  Fast F16 Operation:      TRUE                               
  Wavefront Size:          64(0x40)                           
  Workgroup Max Size:      1024(0x400)                        
  Workgroup Max Size per Dimension:
    x                        1024(0x400)                        
    y                        1024(0x400)                        
    z                        1024(0x400)                        
  Max Waves Per CU:        32(0x20)                           
  Max Work-item Per CU:    2048(0x800)                        
  Grid Max Size:           4294967295(0xffffffff)             
  Grid Max Size per Dimension:
    x                        4294967295(0xffffffff)             
    y                        4294967295(0xffffffff)             
    z                        4294967295(0xffffffff)             
  Max fbarriers/Workgrp:   32                                 
  Packet Processor uCode:: 138                                
  SDMA engine uCode::      19                                 
  IOMMU Support::          None                               
  Pool Info:               
    Pool 1                   
      Segment:                 GLOBAL; FLAGS: COARSE GRAINED      
      Size:                    134217728(0x8000000) KB            
      Allocatable:             TRUE                               
      Alloc Granule:           4KB                                
      Alloc Recommended Granule:2048KB                             
      Alloc Alignment:         4KB                                
      Accessible by all:       FALSE                              
    Pool 2                   
      Segment:                 GLOBAL; FLAGS: EXTENDED FINE GRAINED
      Size:                    134217728(0x8000000) KB            
      Allocatable:             TRUE                               
      Alloc Granule:           4KB                                
      Alloc Recommended Granule:2048KB                             
      Alloc Alignment:         4KB                                
      Accessible by all:       FALSE                              
    Pool 3                   
      Segment:                 GLOBAL; FLAGS: FINE GRAINED        
      Size:                    134217728(0x8000000) KB            
      Allocatable:             TRUE                               
      Alloc Granule:           4KB                                
      Alloc Recommended Granule:2048KB                             
      Alloc Alignment:         4KB                                
      Accessible by all:       FALSE                              
    Pool 4                   
      Segment:                 GROUP                              
      Size:                    64(0x40) KB                        
      Allocatable:             FALSE                              
      Alloc Granule:           0KB                                
      Alloc Recommended Granule:0KB                                
      Alloc Alignment:         0KB                                
      Accessible by all:       FALSE                              
  ISA Info:                
    ISA 1                    
      Name:                    amdgcn-amd-amdhsa--gfx942:sramecc+:xnack+
      Machine Models:          HSA_MACHINE_MODEL_LARGE            
      Profiles:                HSA_PROFILE_BASE                   
      Default Rounding Mode:   NEAR                               
      Default Rounding Mode:   NEAR                               
      Fast f16:                TRUE                               
      Workgroup Max Size:      1024(0x400)                        
      Workgroup Max Size per Dimension:
        x                        1024(0x400)                        
        y                        1024(0x400)                        
        z                        1024(0x400)                        
      Grid Max Size:           4294967295(0xffffffff)             
      Grid Max Size per Dimension:
        x                        4294967295(0xffffffff)             
        y                        4294967295(0xffffffff)             
        z                        4294967295(0xffffffff)             
      FBarrier Max Size:       32                                 
    ISA 2                    
      Name:                    amdgcn-amd-amdhsa--gfx9-4-generic:sramecc+:xnack+
      Machine Models:          HSA_MACHINE_MODEL_LARGE            
      Profiles:                HSA_PROFILE_BASE                   
      Default Rounding Mode:   NEAR                               
      Default Rounding Mode:   NEAR                               
      Fast f16:                TRUE                               
      Workgroup Max Size:      1024(0x400)                        
      Workgroup Max Size per Dimension:
        x                        1024(0x400)                        
        y                        1024(0x400)                        
        z                        1024(0x400)                        
      Grid Max Size:           4294967295(0xffffffff)             
      Grid Max Size per Dimension:
        x                        4294967295(0xffffffff)             
        y                        4294967295(0xffffffff)             
        z                        4294967295(0xffffffff)             
      FBarrier Max Size:       32                                 
*******                  
Agent 6                  
*******                  
  Name:                    gfx942                             
  Uuid:                    GPU-a5b43667c73cac87               
  Marketing Name:          AMD Instinct MI300A                
  Vendor Name:             AMD                                
  Feature:                 KERNEL_DISPATCH                    
  Profile:                 BASE_PROFILE                       
  Float Round Mode:        NEAR                               
  Max Queue Number:        128(0x80)                          
  Queue Min Size:          64(0x40)                           
  Queue Max Size:          131072(0x20000)                    
  Queue Type:              MULTI                              
  Node:                    5                                  
  Device Type:             GPU                                
  Cache Info:              
    L1:                      32(0x20) KB                        
    L2:                      24576(0x6000) KB                   
    L3:                      262144(0x40000) KB                 
  Chip ID:                 29856(0x74a0)                      
  ASIC Revision:           1(0x1)                             
  Cacheline Size:          64(0x40)                           
  Max Clock Freq. (MHz):   2100                               
  BDFID:                   512                                
  Internal Node ID:        5                                  
  Compute Unit:            228                                
  SIMDs per CU:            4                                  
  Shader Engines:          24                                 
  Shader Arrs. per Eng.:   1                                  
  WatchPts on Addr. Ranges:4                                  
  Coherent Host Access:    TRUE                               
  Memory Properties:       APU
  Features:                KERNEL_DISPATCH 
  Fast F16 Operation:      TRUE                               
  Wavefront Size:          64(0x40)                           
  Workgroup Max Size:      1024(0x400)                        
  Workgroup Max Size per Dimension:
    x                        1024(0x400)                        
    y                        1024(0x400)                        
    z                        1024(0x400)                        
  Max Waves Per CU:        32(0x20)                           
  Max Work-item Per CU:    2048(0x800)                        
  Grid Max Size:           4294967295(0xffffffff)             
  Grid Max Size per Dimension:
    x                        4294967295(0xffffffff)             
    y                        4294967295(0xffffffff)             
    z                        4294967295(0xffffffff)             
  Max fbarriers/Workgrp:   32                                 
  Packet Processor uCode:: 138                                
  SDMA engine uCode::      19                                 
  IOMMU Support::          None                               
  Pool Info:               
    Pool 1                   
      Segment:                 GLOBAL; FLAGS: COARSE GRAINED      
      Size:                    134217728(0x8000000) KB            
      Allocatable:             TRUE                               
      Alloc Granule:           4KB                                
      Alloc Recommended Granule:2048KB                             
      Alloc Alignment:         4KB                                
      Accessible by all:       FALSE                              
    Pool 2                   
      Segment:                 GLOBAL; FLAGS: EXTENDED FINE GRAINED
      Size:                    134217728(0x8000000) KB            
      Allocatable:             TRUE                               
      Alloc Granule:           4KB                                
      Alloc Recommended Granule:2048KB                             
      Alloc Alignment:         4KB                                
      Accessible by all:       FALSE                              
    Pool 3                   
      Segment:                 GLOBAL; FLAGS: FINE GRAINED        
      Size:                    134217728(0x8000000) KB            
      Allocatable:             TRUE                               
      Alloc Granule:           4KB                                
      Alloc Recommended Granule:2048KB                             
      Alloc Alignment:         4KB                                
      Accessible by all:       FALSE                              
    Pool 4                   
      Segment:                 GROUP                              
      Size:                    64(0x40) KB                        
      Allocatable:             FALSE                              
      Alloc Granule:           0KB                                
      Alloc Recommended Granule:0KB                                
      Alloc Alignment:         0KB                                
      Accessible by all:       FALSE                              
  ISA Info:                
    ISA 1                    
      Name:                    amdgcn-amd-amdhsa--gfx942:sramecc+:xnack+
      Machine Models:          HSA_MACHINE_MODEL_LARGE            
      Profiles:                HSA_PROFILE_BASE                   
      Default Rounding Mode:   NEAR                               
      Default Rounding Mode:   NEAR                               
      Fast f16:                TRUE                               
      Workgroup Max Size:      1024(0x400)                        
      Workgroup Max Size per Dimension:
        x                        1024(0x400)                        
        y                        1024(0x400)                        
        z                        1024(0x400)                        
      Grid Max Size:           4294967295(0xffffffff)             
      Grid Max Size per Dimension:
        x                        4294967295(0xffffffff)             
        y                        4294967295(0xffffffff)             
        z                        4294967295(0xffffffff)             
      FBarrier Max Size:       32                                 
    ISA 2                    
      Name:                    amdgcn-amd-amdhsa--gfx9-4-generic:sramecc+:xnack+
      Machine Models:          HSA_MACHINE_MODEL_LARGE            
      Profiles:                HSA_PROFILE_BASE                   
      Default Rounding Mode:   NEAR                               
      Default Rounding Mode:   NEAR                               
      Fast f16:                TRUE                               
      Workgroup Max Size:      1024(0x400)                        
      Workgroup Max Size per Dimension:
        x                        1024(0x400)                        
        y                        1024(0x400)                        
        z                        1024(0x400)                        
      Grid Max Size:           4294967295(0xffffffff)             
      Grid Max Size per Dimension:
        x                        4294967295(0xffffffff)             
        y                        4294967295(0xffffffff)             
        z                        4294967295(0xffffffff)             
      FBarrier Max Size:       32                                 
*******                  
Agent 7                  
*******                  
  Name:                    gfx942                             
  Uuid:                    GPU-655f5d89689c7f67               
  Marketing Name:          AMD Instinct MI300A                
  Vendor Name:             AMD                                
  Feature:                 KERNEL_DISPATCH                    
  Profile:                 BASE_PROFILE                       
  Float Round Mode:        NEAR                               
  Max Queue Number:        128(0x80)                          
  Queue Min Size:          64(0x40)                           
  Queue Max Size:          131072(0x20000)                    
  Queue Type:              MULTI                              
  Node:                    6                                  
  Device Type:             GPU                                
  Cache Info:              
    L1:                      32(0x20) KB                        
    L2:                      24576(0x6000) KB                   
    L3:                      262144(0x40000) KB                 
  Chip ID:                 29856(0x74a0)                      
  ASIC Revision:           1(0x1)                             
  Cacheline Size:          64(0x40)                           
  Max Clock Freq. (MHz):   2100                               
  BDFID:                   512                                
  Internal Node ID:        6                                  
  Compute Unit:            228                                
  SIMDs per CU:            4                                  
  Shader Engines:          24                                 
  Shader Arrs. per Eng.:   1                                  
  WatchPts on Addr. Ranges:4                                  
  Coherent Host Access:    TRUE                               
  Memory Properties:       APU
  Features:                KERNEL_DISPATCH 
  Fast F16 Operation:      TRUE                               
  Wavefront Size:          64(0x40)                           
  Workgroup Max Size:      1024(0x400)                        
  Workgroup Max Size per Dimension:
    x                        1024(0x400)                        
    y                        1024(0x400)                        
    z                        1024(0x400)                        
  Max Waves Per CU:        32(0x20)                           
  Max Work-item Per CU:    2048(0x800)                        
  Grid Max Size:           4294967295(0xffffffff)             
  Grid Max Size per Dimension:
    x                        4294967295(0xffffffff)             
    y                        4294967295(0xffffffff)             
    z                        4294967295(0xffffffff)             
  Max fbarriers/Workgrp:   32                                 
  Packet Processor uCode:: 138                                
  SDMA engine uCode::      19                                 
  IOMMU Support::          None                               
  Pool Info:               
    Pool 1                   
      Segment:                 GLOBAL; FLAGS: COARSE GRAINED      
      Size:                    134217728(0x8000000) KB            
      Allocatable:             TRUE                               
      Alloc Granule:           4KB                                
      Alloc Recommended Granule:2048KB                             
      Alloc Alignment:         4KB                                
      Accessible by all:       FALSE                              
    Pool 2                   
      Segment:                 GLOBAL; FLAGS: EXTENDED FINE GRAINED
      Size:                    134217728(0x8000000) KB            
      Allocatable:             TRUE                               
      Alloc Granule:           4KB                                
      Alloc Recommended Granule:2048KB                             
      Alloc Alignment:         4KB                                
      Accessible by all:       FALSE                              
    Pool 3                   
      Segment:                 GLOBAL; FLAGS: FINE GRAINED        
      Size:                    134217728(0x8000000) KB            
      Allocatable:             TRUE                               
      Alloc Granule:           4KB                                
      Alloc Recommended Granule:2048KB                             
      Alloc Alignment:         4KB                                
      Accessible by all:       FALSE                              
    Pool 4                   
      Segment:                 GROUP                              
      Size:                    64(0x40) KB                        
      Allocatable:             FALSE                              
      Alloc Granule:           0KB                                
      Alloc Recommended Granule:0KB                                
      Alloc Alignment:         0KB                                
      Accessible by all:       FALSE                              
  ISA Info:                
    ISA 1                    
      Name:                    amdgcn-amd-amdhsa--gfx942:sramecc+:xnack+
      Machine Models:          HSA_MACHINE_MODEL_LARGE            
      Profiles:                HSA_PROFILE_BASE                   
      Default Rounding Mode:   NEAR                               
      Default Rounding Mode:   NEAR                               
      Fast f16:                TRUE                               
      Workgroup Max Size:      1024(0x400)                        
      Workgroup Max Size per Dimension:
        x                        1024(0x400)                        
        y                        1024(0x400)                        
        z                        1024(0x400)                        
      Grid Max Size:           4294967295(0xffffffff)             
      Grid Max Size per Dimension:
        x                        4294967295(0xffffffff)             
        y                        4294967295(0xffffffff)             
        z                        4294967295(0xffffffff)             
      FBarrier Max Size:       32                                 
    ISA 2                    
      Name:                    amdgcn-amd-amdhsa--gfx9-4-generic:sramecc+:xnack+
      Machine Models:          HSA_MACHINE_MODEL_LARGE            
      Profiles:                HSA_PROFILE_BASE                   
      Default Rounding Mode:   NEAR                               
      Default Rounding Mode:   NEAR                               
      Fast f16:                TRUE                               
      Workgroup Max Size:      1024(0x400)                        
      Workgroup Max Size per Dimension:
        x                        1024(0x400)                        
        y                        1024(0x400)                        
        z                        1024(0x400)                        
      Grid Max Size:           4294967295(0xffffffff)             
      Grid Max Size per Dimension:
        x                        4294967295(0xffffffff)             
        y                        4294967295(0xffffffff)             
        z                        4294967295(0xffffffff)             
      FBarrier Max Size:       32                                 
*******                  
Agent 8                  
*******                  
  Name:                    gfx942                             
  Uuid:                    GPU-5f8fb5cf1f014be1               
  Marketing Name:          AMD Instinct MI300A                
  Vendor Name:             AMD                                
  Feature:                 KERNEL_DISPATCH                    
  Profile:                 BASE_PROFILE                       
  Float Round Mode:        NEAR                               
  Max Queue Number:        128(0x80)                          
  Queue Min Size:          64(0x40)                           
  Queue Max Size:          131072(0x20000)                    
  Queue Type:              MULTI                              
  Node:                    7                                  
  Device Type:             GPU                                
  Cache Info:              
    L1:                      32(0x20) KB                        
    L2:                      24576(0x6000) KB                   
    L3:                      262144(0x40000) KB                 
  Chip ID:                 29856(0x74a0)                      
  ASIC Revision:           1(0x1)                             
  Cacheline Size:          64(0x40)                           
  Max Clock Freq. (MHz):   2100                               
  BDFID:                   512                                
  Internal Node ID:        7                                  
  Compute Unit:            228                                
  SIMDs per CU:            4                                  
  Shader Engines:          24                                 
  Shader Arrs. per Eng.:   1                                  
  WatchPts on Addr. Ranges:4                                  
  Coherent Host Access:    TRUE                               
  Memory Properties:       APU
  Features:                KERNEL_DISPATCH 
  Fast F16 Operation:      TRUE                               
  Wavefront Size:          64(0x40)                           
  Workgroup Max Size:      1024(0x400)                        
  Workgroup Max Size per Dimension:
    x                        1024(0x400)                        
    y                        1024(0x400)                        
    z                        1024(0x400)                        
  Max Waves Per CU:        32(0x20)                           
  Max Work-item Per CU:    2048(0x800)                        
  Grid Max Size:           4294967295(0xffffffff)             
  Grid Max Size per Dimension:
    x                        4294967295(0xffffffff)             
    y                        4294967295(0xffffffff)             
    z                        4294967295(0xffffffff)             
  Max fbarriers/Workgrp:   32                                 
  Packet Processor uCode:: 138                                
  SDMA engine uCode::      19                                 
  IOMMU Support::          None                               
  Pool Info:               
    Pool 1                   
      Segment:                 GLOBAL; FLAGS: COARSE GRAINED      
      Size:                    134217728(0x8000000) KB            
      Allocatable:             TRUE                               
      Alloc Granule:           4KB                                
      Alloc Recommended Granule:2048KB                             
      Alloc Alignment:         4KB                                
      Accessible by all:       FALSE                              
    Pool 2                   
      Segment:                 GLOBAL; FLAGS: EXTENDED FINE GRAINED
      Size:                    134217728(0x8000000) KB            
      Allocatable:             TRUE                               
      Alloc Granule:           4KB                                
      Alloc Recommended Granule:2048KB                             
      Alloc Alignment:         4KB                                
      Accessible by all:       FALSE                              
    Pool 3                   
      Segment:                 GLOBAL; FLAGS: FINE GRAINED        
      Size:                    134217728(0x8000000) KB            
      Allocatable:             TRUE                               
      Alloc Granule:           4KB                                
      Alloc Recommended Granule:2048KB                             
      Alloc Alignment:         4KB                                
      Accessible by all:       FALSE                              
    Pool 4                   
      Segment:                 GROUP                              
      Size:                    64(0x40) KB                        
      Allocatable:             FALSE                              
      Alloc Granule:           0KB                                
      Alloc Recommended Granule:0KB                                
      Alloc Alignment:         0KB                                
      Accessible by all:       FALSE                              
  ISA Info:                
    ISA 1                    
      Name:                    amdgcn-amd-amdhsa--gfx942:sramecc+:xnack+
      Machine Models:          HSA_MACHINE_MODEL_LARGE            
      Profiles:                HSA_PROFILE_BASE                   
      Default Rounding Mode:   NEAR                               
      Default Rounding Mode:   NEAR                               
      Fast f16:                TRUE                               
      Workgroup Max Size:      1024(0x400)                        
      Workgroup Max Size per Dimension:
        x                        1024(0x400)                        
        y                        1024(0x400)                        
        z                        1024(0x400)                        
      Grid Max Size:           4294967295(0xffffffff)             
      Grid Max Size per Dimension:
        x                        4294967295(0xffffffff)             
        y                        4294967295(0xffffffff)             
        z                        4294967295(0xffffffff)             
      FBarrier Max Size:       32                                 
    ISA 2                    
      Name:                    amdgcn-amd-amdhsa--gfx9-4-generic:sramecc+:xnack+
      Machine Models:          HSA_MACHINE_MODEL_LARGE            
      Profiles:                HSA_PROFILE_BASE                   
      Default Rounding Mode:   NEAR                               
      Default Rounding Mode:   NEAR                               
      Fast f16:                TRUE                               
      Workgroup Max Size:      1024(0x400)                        
      Workgroup Max Size per Dimension:
        x                        1024(0x400)                        
        y                        1024(0x400)                        
        z                        1024(0x400)                        
      Grid Max Size:           4294967295(0xffffffff)             
      Grid Max Size per Dimension:
        x                        4294967295(0xffffffff)             
        y                        4294967295(0xffffffff)             
        z                        4294967295(0xffffffff)             
      FBarrier Max Size:       32                                 
*** Done *** 

### Additional Information

no specific message regarding the location of memory in dmesg, nor with AMD_LOG_LEVEL

[etiennemlb]

Up please

[ppanchad-amd]

Hi @cboillot. Internal ticket has been created to investigate this issue. Thanks!

[tcgu-amd]

Hi @cboillot, sorry to hear that you are having trouble with ROCm in MI300A, and thank you for reaching out. What you are describing is definitely not normal, since hipMalloc targets specific device and there's no mechanism for it to "borrow" from other device's memory as far as I am aware. Do you have a reproducer with HIP? Pytorch does a lot of stuff under the hood so it is hard to pin down the issue. Thanks!

[cboillot]

Hello,
Thanks for the response, here is a HIP reproducer to show that hip "borrow" memory from another device's memory :

$ ./repro.sh 
start filling the pagecache
reuse file
Memory Info:
Node       Used GiB   Cached GiB Free GiB  
0          114        103        9
1          18         15         107       
2          1          0          123       
3          1          0          123       
Allocated 110 tensors of size 1 GiB on GPU 0.
Memory Info:
Node       Used GiB   Cached GiB Free GiB  
0          116        86         8        <-- The pagecache still use 86 GiB (+ overhead) => no more than 30GiB comes from GPU 0 memory
1          108        15         17      <-- Approximately 90 GiB of memory is in use on GPU 1, despite allocations being made only on GPU 0.
2          2          0          123       
3          1          0          123 

As you can see despite the hipSetDevice(0); due to the fact that numa node 0 has been filled by numactl --cpubind=0 --membind=0 cat ./zero.bin > /dev/null
the 110 GiB of memory allocated on GPU 0 mostly come from numa node 1 :

reproducer files :
repro.cxx

#include <hip/hip_runtime.h>
#include <iostream>
#include <vector>
#include <cstdlib> // For system()

void print_memory_info() {
    // Execute the command to print memory information
    std::cout << "Memory Info:\n";
    system("./numa_info.sh");
}

int main() {
    const int num_tensors = 110;
    const size_t tensor_size = 1 * 1024 * 1024 * 1024; // Size in bytes (1 GiB)

    // Set the device to GPU 0
    hipError_t setDeviceErr = hipSetDevice(0);
    if (setDeviceErr != hipSuccess) {
        std::cerr << "Error setting device: " << hipGetErrorString(setDeviceErr) << std::endl;
        return -1;
    }

    // Print memory info before allocation
    print_memory_info();

    // Vector to hold device pointers
    std::vector<float*> device_tensors(num_tensors);

    // Allocate memory on GPU
    for (int i = 0; i < num_tensors; ++i) {
        hipError_t err = hipMalloc(&device_tensors[i], tensor_size);
        if (err != hipSuccess) {
            std::cerr << "Error allocating tensor " << i << ": " << hipGetErrorString(err) << std::endl;
            return -1;
        }
    }

    std::cout << "Allocated " << num_tensors << " tensors of size 1 GiB on GPU 0." << std::endl;

    // Print memory info after allocation
    print_memory_info();

    // Free the allocated memory
    for (int i = 0; i < num_tensors; ++i) {
        hipError_t err = hipFree(device_tensors[i]);
        if (err != hipSuccess) {
            std::cerr << "Error freeing tensor " << i << ": " << hipGetErrorString(err) << std::endl;
        }
    }

    return 0;
}

repro.sh

#!/bin/bash
set -eu

hipcc repro.cxx -o repro
echo "start filling the pagecache"
if [ -f ./zero.bin ]; then
        echo "reuse file"
        numactl --cpubind=0 --membind=0 cat ./zero.bin > /dev/null
else 
        echo "build a new zero file"
        numactl --cpubind=0 --membind=0 dd if=/dev/zero of=./zero.bin bs=1G count=150 status=progress
        sync
fi
./repro

numa_info.sh

grep -hi 'MemFree\|FilePages\|MemUsed' /sys/devices/system/node/node*/meminfo | awk '
BEGIN {
    printf "%-10s %-10s %-10s %-10s\n", "Node", "Used GiB", "Cached GiB", "Free GiB"
}
{
    if (/Node/) {
        node = $2
    }
    if (/MemUsed:/) {
        mem_info[node]["used"] = $4/1024/1024
    } else if (/FilePages:/) {
        mem_info[node]["cached"] = $4/1024/1024
    } else if (/MemFree:/) {
        mem_info[node]["free"] = $4/1024/1024
    }
}
END {
    for (n in mem_info) {
        printf "%-10s %-10d %-10d %-10d\n", n, mem_info[n]["used"], mem_info[n]["cached"], mem_info[n]["free"]
    }
}'

[cboillot]

Hello @tcgu-amd,

Have you had the time to reproduce and investigate the issue?

Below is a second reproducer to check that the allocation fails when the GPU memory is exhausted :

hipcc repro_v2.cxx -lnuma -o repro_v2

./repro_v2 
Memory Info after numa alloc on numa 0 and 1
Node       Used GiB   Cached GiB Free GiB  
0          65         0          58        
1          66         0          59        
2          1          0          123       
3          1          0          123       
Allocated 93 tensors of size 1 GiB on GPU 0.
Memory Info after NUMA + GPU allocation
Node       Used GiB   Cached GiB Free GiB  
0          119        0          5         
1          106        0          18        
2          1          0          123       
3          1          0          123       
Memory Info after  NUMA  free and before GPU free
Node       Used GiB   Cached GiB Free GiB  
0          55         0          69        
1          42         0          83        
2          1          0          123       
3          1          0          123  

repro_v2.cxx
#include <hip/hip_runtime.h>
#include <iostream>
#include <vector>
#include <cstdlib>
#include <numa.h>
#include <numaif.h>

void print_memory_info(std::string when) {
    // Execute the command to print memory information
    std::cout << "Memory Info " << when << "\n";
    system("./numa_info.sh");
}

int main() {
    const int num_tensors = 93;
    const size_t tensor_size = 1 * 1024 * 1024 * 1024; // Size in bytes (1 GiB)
    const size_t numa_alloc_size = 64ULL * 1024 * 1024 * 1024;

    // Allocate 64 GiB on NUMA node 0
    void *numa_ptr = numa_alloc_onnode(numa_alloc_size, 0);
    if (numa_ptr == NULL) {
        fprintf(stderr, "Failed to allocate memory on NUMA node %d.\n", 0);
        return EXIT_FAILURE;
    }

    // force pagefault
    memset(numa_ptr, 0, numa_alloc_size);

    // Allocate 64 GiB on NUMA node 1
    void *numa_ptr_1 = numa_alloc_onnode(numa_alloc_size, 1);
    if (numa_ptr_1 == NULL) {
        fprintf(stderr, "Failed to allocate memory on NUMA node %d.\n", 1);
        return EXIT_FAILURE;
    }

    // force pagefault
    memset(numa_ptr_1, 0, numa_alloc_size);

    // Set the device to GPU 0
    hipError_t setDeviceErr = hipSetDevice(0);
    if (setDeviceErr != hipSuccess) {
        std::cerr << "Error setting device: " << hipGetErrorString(setDeviceErr) << std::endl;
        return -1;
    }

    // Print memory info before allocation
    print_memory_info("after numa alloc on numa 0 and 1");

    // Vector to hold device pointers
    std::vector<float*> device_tensors(num_tensors);

    // Allocate memory on GPU
    for (int i = 0; i < num_tensors; ++i) {
        hipError_t err = hipMalloc(&device_tensors[i], tensor_size);
        if (err != hipSuccess) {
            std::cerr << "Error allocating tensor " << i << ": " << hipGetErrorString(err) << std::endl;
            device_tensors[i] = NULL;
        }
    }

    std::cout << "Allocated " << num_tensors << " tensors of size 1 GiB on GPU 0." << std::endl;

    // Print memory info after allocation
    print_memory_info("after NUMA + GPU allocation");

    numa_free(numa_ptr, numa_alloc_size);
    numa_free(numa_ptr_1, numa_alloc_size);

    print_memory_info("after  NUMA  free and before GPU free"); 
    // Free the allocated memory
    for (int i = 0; i < num_tensors; ++i) {
        if (device_tensors[i]) {
            hipError_t err = hipFree(device_tensors[i]);
            if (err != hipSuccess) {
                std::cerr << "Error freeing tensor " << i << ": " << hipGetErrorString(err) << std::endl;
            }
        }
    }
 

    return 0;
}

[tcgu-amd]

@cboillot, yes, I was able to reproduce your issue. We are currently in the process of analyzing the root cause. Thank you!

[etiennemlb]

Great, we are all impatiently waiting for your answer. Have a nice day.

…

On 10/07/2025 20:18, Tim Gu wrote: *tcgu-amd* left a comment (ROCm/amdgpu#186) <https://github.com/ROCm/ amdgpu/issues/186#issuecomment-3058462487> @cboillot <https://github.com/cboillot>, yes, I was able to reproduce your issue. We are currently in the process of analyzing the root cause. Thank you! — Reply to this email directly, view it on GitHub <https://github.com/ ROCm/amdgpu#186#issuecomment-3058462487>, or unsubscribe <https://github.com/notifications/unsubscribe-auth/ ARHSPD3PMYTHASMBDEPZ5N33H2U6PAVCNFSM6AAAAAB5NRON72VHI2DSMVQWIX3LMV43OSLTON2WKQ3PNVWWK3TUHMZTANJYGQ3DENBYG4>. You are receiving this because you are subscribed to this thread.Message ID: ***@***.***>

[tcgu-amd]

Hey @etiennemlb, there's actually a patch that should address this issue. Can you please check driver source code under /usr/src/amdgpu-xx.xx to see if this patch exist?

commit 9c27f61
Author: Rajneesh Bhardwaj <rajneesh.bhardwaj@amd.com>
Date: Mon Jul 8 12:06:36 2024 -0400

drm/ttm: Allow direct reclaim to allocate local memory v2

Limiting the allocation of higher order pages to the closest NUMA node
and enabling direct memory reclaim provides not only failsafe against
situations when memory becomes too much fragmented and the allocator is
not able to satisfy the request from the local node but falls back to
remote pages (HUGEPAGE) but also offers performance improvement.
Accessing remote pages suffers due to bandwidth limitations and could be
avoided if memory becomes defragmented and in most cases without using
manual compaction. (/proc/sys/vm/compact_memory)

Note: On certain distros such as RHEL, the proactive compaction is
Note: On certain distros such as RHEL, the proactive compaction is
disabled. (https://tinyurl.com/4f32f7rs)

v2 (chk): drop __GFP_RECLAIM since that is already set by GFP_USER

Cc: Dave Airlie <[airlied@redhat.com](mailto:airlied@redhat.com)>
Cc: Vlastimil Babka <[vbabka@suse.cz](mailto:vbabka@suse.cz)>
Cc: Daniel Vetter <[daniel.vetter@ffwll.ch](mailto:daniel.vetter@ffwll.ch)>
Reviewed-by: Christian König <[christian.koenig@amd.com](mailto:christian.koenig@amd.com)>
Signed-off-by: Rajneesh Bhardwaj <[rajneesh.bhardwaj@amd.com](mailto:rajneesh.bhardwaj@amd.com)>
Link: https://patchwork.freedesktop.org/patch/msgid/20240708160636.1147308-1-rajneesh.bhardwaj@amd.com
Signed-off-by: Christian König <[christian.koenig@amd.com](mailto:christian.koenig@amd.com)>

diff --git a/drivers/gpu/drm/ttm/ttm_pool.c b/drivers/gpu/drm/ttm/ttm_pool.c
index 6e1fd6985ffc..8504dbe19c1a 100644
--- a/drivers/gpu/drm/ttm/ttm_pool.c
+++ b/drivers/gpu/drm/ttm/ttm_pool.c
@@ -91,7 +91,7 @@ static struct page *ttm_pool_alloc_page(struct ttm_pool *pool, gfp_t gfp_flags,
*/
if (order)
gfp_flags |= __GFP_NOMEMALLOC | __GFP_NORETRY | __GFP_NOWARN |

•                   __GFP_KSWAPD_RECLAIM;
  

•                   __GFP_THISNODE;
  
    if (!pool->use_dma_alloc) {
            p = alloc_pages_node(pool->nid, gfp_flags, order);
  

[etiennemlb]

Ok I just checked if we had this fix, and it turns out we do not. I'll see if we can try this patch.

[GeirHPE]

Which version of AMD ROCm has or will have the patch fix included in the release? Thanks.

[tcgu-amd]

Which version of AMD ROCm has or will have the patch fix included in the release? Thanks.

This patched was merged a while back and should be available with amdgpu driver released recently (the patch was merged at least a year ago).

[tcgu-amd]

Hi @etiennemlb, a bit of an update from our side again. So, there are quite a few things going on. HIP will try to allocate memory on each device as specified, however, the kernel is free to override this in two ways. The first behavior is to try allocating on a different NUMA device when allocation on one NUMA fails, which is actually standard Linux kernel behavior. As far as I understand, the patch above tries to override this behavior. The second way is to evict memory from one device to another, so the requested memory allocation falls on the same device, but it may appear that memory is allocated on a different one. Because the issue is really just standard Linux memory management behavior, memory already allocated by HIP can also be evicted by different, potentially non-HIP processes as well.

Another caveat is that HIP is written such that it gauges available device memory on APUs based on the assumption that memory on the device can be evicted. The issue is that with MI300A being an APU, the evicted memory will be evicted to yet another device, so under really high memory pressure with all devices utilized, eventually the assumption will fail. That's why typically device page_limit is configured to be 3/4 of the available system memory to avoid this.

tl;dr, so far, the reported behavior actually looks expected. It works that way due to the special nature of APUs where device/system memory is shared. This behavior should be able to be controlled by avoiding situations that leads to extremely high memory pressure. In the meantime, there are still on-going discussions regarding if this should be the optimal behavior of HIP on APUs, but there are a lot of pros and cons for each option and the current one looks the most compliant with Linux kernel specifications.

[cboillot]

Hello @tcgu-amd,

I confirm that the driver update with the patch seems to effectively resolve the encountered issue. Once the update is properly deployed on our side, we will conduct another round of testing to validate the behavior at the limits (from what I read, __GFP_THISNODE does indeed seem to be the right solution).