[Issue]: GPU acceleration crashes (Cannot find Symbol with name)

[jakobwinkler]

Problem Description

Hi,

I'm having issues getting the GPU acceleration to work. Using an RX 7900 XTX on Fedora Workstation 40, Kernel 6.9.9-200.fc40.x86_64.

I built the library from both master-rocm and v2.0.3-rocm using the instructions in README-ROCm.md (so mainly running cmake -DUSE_HIP=1 ../ and make) with the same result.

Using demo/CLI/binary_classification as a reproduction case, it runs normally in its default configuration, but adding device = "gpu" to mushroom.conf leads to the following issue, other reproduction cases show similar behavior, but with different symbols depending on the usage:

:0:/builddir/build/BUILD/clr-rocm-6.0.2/hipamd/src/hip_global.cpp:114 : 6326756172 us: [pid:35900 tid:0x7fc3aa1142c0] Cannot find Symbol with name: _ZN7xgboost4tree20EvaluateSplitsKernelILi64EEEvjNS_6common4SpanIKNS0_19EvaluateSplitInputsELm18446744073709551615EEENS0_25EvaluateSplitSharedInputsENS3_IjLm18446744073709551615EEENS0_13TreeEvaluator14SplitEvaluatorINS0_16GPUTrainingParamEEENS3_INS0_20DeviceSplitCandidateELm18446744073709551615EEE 

./runexp.sh: line 10: 35900 Aborted                 (core dumped) $XGBOOST mushroom.conf

Unfortunately I'm unfamiliar with the inner workings of GPU acceleration, but what I've verified is that the .so file contains the symbol.

$ nm -gD ../../../lib/libxgboost.so | grep _ZN7xgboost4tree20EvaluateSplitsKernelILi64EEEvjNS_6common4SpanIKNS0_19EvaluateSplitInputsELm18446744073709551615EEENS0_25EvaluateSplitSharedInputsENS3_IjLm18446744073709551615EEENS0_13TreeEvaluator14SplitEvaluatorINS0_16GPUTrainingParamEEENS3_INS0_20DeviceSplitCandidateELm18446744073709551615EEE
0000000002186038 V _ZN7xgboost4tree20EvaluateSplitsKernelILi64EEEvjNS_6common4SpanIKNS0_19EvaluateSplitInputsELm18446744073709551615EEENS0_25EvaluateSplitSharedInputsENS3_IjLm18446744073709551615EEENS0_13TreeEvaluator14SplitEvaluatorINS0_16GPUTrainingParamEEENS3_INS0_20DeviceSplitCandidateELm18446744073709551615EEE

In general, calls to the library seem to work (see attached logs for ltrace -e "hip*" and a run with AMD_LOG_LEVEL=3). Additionally I attached the output of rocminfo and a gdb backtrace of the crash.

Any pointers would be very welcome.
gdb.log
loglevel3.log
ltrace.log
rocminfo.log

Operating System

Fedora Linux 40 (Workstation Edition) x86_64

CPU

13th Gen Intel(R) Core(TM) i5-13600KF

GPU

AMD Radeon RX 7900 XTX

ROCm Version

ROCm 6.0.0

ROCm Component

No response

Steps to Reproduce

No response

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

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

==========               
HSA Agents               
==========               
*******                  
Agent 1                  
*******                  
  Name:                    13th Gen Intel(R) Core(TM) i5-13600KF
  Uuid:                    CPU-XX                             
  Marketing Name:          13th Gen Intel(R) Core(TM) i5-13600KF
  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:                      49152(0xc000) KB                   
  Chip ID:                 0(0x0)                             
  ASIC Revision:           0(0x0)                             
  Cacheline Size:          64(0x40)                           
  Max Clock Freq. (MHz):   5100                               
  BDFID:                   0                                  
  Internal Node ID:        0                                  
  Compute Unit:            20                                 
  SIMDs per CU:            0                                  
  Shader Engines:          0                                  
  Shader Arrs. per Eng.:   0                                  
  WatchPts on Addr. Ranges:1                                  
  Features:                None
  Pool Info:               
    Pool 1                   
      Segment:                 GLOBAL; FLAGS: FINE GRAINED        
      Size:                    32674140(0x1f2915c) KB             
      Allocatable:             TRUE                               
      Alloc Granule:           4KB                                
      Alloc Alignment:         4KB                                
      Accessible by all:       TRUE                               
    Pool 2                   
      Segment:                 GLOBAL; FLAGS: KERNARG, FINE GRAINED
      Size:                    32674140(0x1f2915c) KB             
      Allocatable:             TRUE                               
      Alloc Granule:           4KB                                
      Alloc Alignment:         4KB                                
      Accessible by all:       TRUE                               
    Pool 3                   
      Segment:                 GLOBAL; FLAGS: COARSE GRAINED      
      Size:                    32674140(0x1f2915c) KB             
      Allocatable:             TRUE                               
      Alloc Granule:           4KB                                
      Alloc Alignment:         4KB                                
      Accessible by all:       TRUE                               
  ISA Info:                
*******                  
Agent 2                  
*******                  
  Name:                    gfx1100                            
  Uuid:                    GPU-2764dec95c0d1c37               
  Marketing Name:          AMD Radeon RX 7900 XTX             
  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:                    1                                  
  Device Type:             GPU                                
  Cache Info:              
    L1:                      32(0x20) KB                        
    L2:                      6144(0x1800) KB                    
    L3:                      98304(0x18000) KB                  
  Chip ID:                 29772(0x744c)                      
  ASIC Revision:           0(0x0)                             
  Cacheline Size:          64(0x40)                           
  Max Clock Freq. (MHz):   2526                               
  BDFID:                   768                                
  Internal Node ID:        1                                  
  Compute Unit:            96                                 
  SIMDs per CU:            2                                  
  Shader Engines:          6                                  
  Shader Arrs. per Eng.:   2                                  
  WatchPts on Addr. Ranges:4                                  
  Coherent Host Access:    FALSE                              
  Features:                KERNEL_DISPATCH 
  Fast F16 Operation:      TRUE                               
  Wavefront Size:          32(0x20)                           
  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:    1024(0x400)                        
  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:: 202                                
  SDMA engine uCode::      21                                 
  IOMMU Support::          None                               
  Pool Info:               
    Pool 1                   
      Segment:                 GLOBAL; FLAGS: COARSE GRAINED      
      Size:                    25149440(0x17fc000) KB             
      Allocatable:             TRUE                               
      Alloc Granule:           4KB                                
      Alloc Alignment:         4KB                                
      Accessible by all:       FALSE                              
    Pool 2                   
      Segment:                 GLOBAL; FLAGS: EXTENDED FINE GRAINED
      Size:                    25149440(0x17fc000) KB             
      Allocatable:             TRUE                               
      Alloc Granule:           4KB                                
      Alloc Alignment:         4KB                                
      Accessible by all:       FALSE                              
    Pool 3                   
      Segment:                 GROUP                              
      Size:                    64(0x40) KB                        
      Allocatable:             FALSE                              
      Alloc Granule:           0KB                                
      Alloc Alignment:         0KB                                
      Accessible by all:       FALSE                              
  ISA Info:                
    ISA 1                    
      Name:                    amdgcn-amd-amdhsa--gfx1100         
      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 response

[hliuca]

Sorry @jakobwinkler, I just noticed the issue. I don't have access to Radeon GPUs, which have different architecture, like wave32, while data center GPUs have wave64.

src/tree/gpu_hist/evaluate_splits.cu

#if defined(XGBOOST_USE_CUDA)
#define WARP_SIZE 32
#elif defined(XGBOOST_USE_HIP)
#include <hip/hip_cooperative_groups.h>

#ifdef __AMDGCN_WAVEFRONT_SIZE
#undef WAVEFRONT_SIZE
#define WAVEFRONT_SIZE __AMDGCN_WAVEFRONT_SIZE
#endif

#define WARP_SIZE WAVEFRONT_SIZE
#endif

I also see the OS is Fedora. One workaround is to change WAVEFRONT_SIZE to 32. The newer version ROCm has better support for Radeon cards.