HIP Compilation error on Nvidia hardware

[fluidnumerics-joe]

This issue was first reported in the hipfort repository : ROCm/hipfort#36 , but seems to be coming from an issue related to HIP.

System Info
Operating System : Ubuntu 20.04
GPU : Nvidia GTX 1070
CUDA Toolkit 11.0 installed
HIP Installed from ROCm 3.7.0

Source Code
https://github.com/ROCmSoftwarePlatform/hipfort/blob/master/test/vecadd/hip_implementation.cpp

Error message
To recreate the error, Compile the above listed hip_implementation.cpp with hipcc on NVCC platform.
I have found the error only occurs on Nvidia hardware. I've compiled using the command

HIP_PLATFORM=nvcc /opt/rocm/bin/hipcc "--gpu-architecture=sm_61" -c hip_implementation.cpp -o /home/joe/apps/hipfort/test/vecadd/hip_implementation.o

The following output ensues :

$ HIP_PLATFORM=nvcc /opt/rocm/bin/hipcc "--gpu-architecture=sm_61" -c hip_implementation.cpp -o /home/joe/apps/hipfort/test/vecadd/hip_implementation.o
hip_implementation.cpp: In function ‘void vector_add(double*, double*, double*, int)’:
hip_implementation.cpp:6:18: error: ‘blockIdx’ was not declared in this scope
    6 |   size_t index = blockIdx.x * blockDim.x + threadIdx.x;
      |                  ^~~~~~~~
hip_implementation.cpp:6:31: error: ‘blockDim’ was not declared in this scope
    6 |   size_t index = blockIdx.x * blockDim.x + threadIdx.x;
      |                               ^~~~~~~~
hip_implementation.cpp:6:44: error: ‘threadIdx’ was not declared in this scope
    6 |   size_t index = blockIdx.x * blockDim.x + threadIdx.x;
      |                                            ^~~~~~~~~
hip_implementation.cpp:7:32: error: ‘gridDim’ was not declared in this scope
    7 |   size_t stride = blockDim.x * gridDim.x;
      |                                ^~~~~~~
In file included from /opt/rocm-3.7.0/hip/include/hip/hip_runtime.h:58,
                 from hip_implementation.cpp:1:
hip_implementation.cpp: In function ‘void launch(double**, double**, double**, int)’:
hip_implementation.cpp:19:5: error: expected primary-expression before ‘<’ token
   19 |     hipLaunchKernelGGL((vector_add), dim3(320), dim3(256), 0, 0, *dout, *da, *db, N);
      |     ^~~~~~~~~~~~~~~~~~
hip_implementation.cpp:19:5: error: expected primary-expression before ‘>’ token
   19 |     hipLaunchKernelGGL((vector_add), dim3(320), dim3(256), 0, 0, *dout, *da, *db, N);
      |     ^~~~~~~~~~~~~~~~~~
make: *** [Makefile:19: /home/joe/apps/hipfort/test/vecadd/hip_implementation.o] Error 1

Outside of the undefined grid/thread/block errors, hipcc seems to now have compilation issues

hip_implementation.cpp:19:5: error: expected primary-expression before ‘<’ token

This was not an issue with ROCm 3.3.0 installation of HIP.

[fluidnumerics-joe]

Looking at include/hip/nvcc_detail/hip_runtime.h shows that __CUDACC__ must be defined in order for hipBlockIDx, etc. to be defined in this header.

77 #ifdef __CUDACC__
 78 
 79 
 80 #define hipThreadIdx_x threadIdx.x
 81 #define hipThreadIdx_y threadIdx.y
 82 #define hipThreadIdx_z threadIdx.z
 83 
 84 #define hipBlockIdx_x blockIdx.x
 85 #define hipBlockIdx_y blockIdx.y
 86 #define hipBlockIdx_z blockIdx.z
 87 
 88 #define hipBlockDim_x blockDim.x
 89 #define hipBlockDim_y blockDim.y
 90 #define hipBlockDim_z blockDim.z
 91 
 92 #define hipGridDim_x gridDim.x
 93 #define hipGridDim_y gridDim.y
 94 #define hipGridDim_z gridDim.z
 95 
 96 #define HIP_SYMBOL(X) &X

Verbose compilation with hipcc shows that __CUDACC__ is not defined.

$ HIP_PLATFORM=nvcc /opt/rocm/bin/hipcc -v "--gpu-architecture=sm_61" -c hip_implementation.cpp -o /home/joe/apps/hipfort/test/vecadd/hip_implementation.o
#$ _NVVM_BRANCH_=nvvm
#$ _SPACE_= 
#$ _CUDART_=cudart
#$ _HERE_=/usr/local/cuda/bin
#$ _THERE_=/usr/local/cuda/bin
#$ _TARGET_SIZE_=
#$ _TARGET_DIR_=
#$ _TARGET_DIR_=targets/x86_64-linux
#$ TOP=/usr/local/cuda/bin/..
#$ NVVMIR_LIBRARY_DIR=/usr/local/cuda/bin/../nvvm/libdevice
#$ LD_LIBRARY_PATH=/usr/local/cuda/bin/../lib:/opt/rocm//lib:/opt/json-fortran//lib:/opt/hipfort//lib:/opt/feqparse//lib
#$ PATH=/usr/local/cuda/bin/../nvvm/bin:/usr/local/cuda/bin:/opt/rocm//bin:/opt/rocm//include:/opt/json-fortran//lib:/opt/hipfort//bin:/opt/hipfort//include:/opt/feqparse//bin:/opt/feqparse//include:/usr/local/Modules/bin:/usr/local/sbin:/usr/local/bin:/usr/sbin:/usr/bin:/sbin:/bin:/usr/games:/usr/local/games:/snap/bin:/home/joe/apps/paraview/bin
#$ INCLUDES="-I/usr/local/cuda/bin/../targets/x86_64-linux/include"  
#$ LIBRARIES=  "-L/usr/local/cuda/bin/../targets/x86_64-linux/lib/stubs" "-L/usr/local/cuda/bin/../targets/x86_64-linux/lib"
#$ CUDAFE_FLAGS=
#$ PTXAS_FLAGS=
#$ gcc -c -x c++ -D__NVCC__  "-I/usr/local/cuda/bin/../targets/x86_64-linux/include"   -isystem "/usr/local/cuda/include" -isystem "/opt/rocm-3.7.0/hip/include"  -D "__HIP_ROCclr__" -D__CUDACC_VER_MAJOR__=11 -D__CUDACC_VER_MINOR__=0 -D__CUDACC_VER_BUILD__=221 -m64 "hip_implementation.cpp" -o "/home/joe/apps/hipfort/test/vecadd/hip_implementation.o" 

[fluidnumerics-joe]

Manually setting

export HIP_COMPILER=/usr/local/cuda/bin/nvcc

resolves the compilation issue.

Additionally, I found the following modification to hipcc also results in successful compilation (with HIP_COMPILER unset in my environment )

323 } elsif ($HIP_PLATFORM eq "nvcc") {
324     $CUDA_PATH=$ENV{'CUDA_PATH'} // '/usr/local/cuda';
325     $HIP_INCLUDE_PATH = "$HIP_PATH/include";
326     $HIP_COMPILER="$CUDA_PATH/bin/nvcc"; # Set the HIP_COMPILER
327     if ($verbose & 0x2) {
328         print ("CUDA_PATH=$CUDA_PATH\n");
329     }
330 
331     $HIPCC="$CUDA_PATH/bin/nvcc";
332     $HIPCXXFLAGS .= " -Wno-deprecated-gpu-targets ";
333     $HIPCXXFLAGS .= " -isystem $CUDA_PATH/include";
334     $HIPCFLAGS .= " -isystem $CUDA_PATH/include";
335 
336     $HIPLDFLAGS = " -Wno-deprecated-gpu-targets -lcuda -lcudart -L$CUDA_PATH/lib64";
337 } else {
338     printf ("error: unknown HIP_PLATFORM = '$HIP_PLATFORM'");
339     printf ("       or HIP_COMPILER = '$HIP_COMPILER'");
340     exit (-1);
341 }

[satyanveshd]

HIP support for nvcc has been refactored already in the internal dev branch and this issue is addressed as part of it. This change should propagate to the development branch soon. Thanks.

[fluidnumerics-joe]

@satyanveshd thanks for the update! Is there any CI testing that is done on Nvidia platforms to ensure that major releases are not broken ? If not, I'd be interested in joining a maintainers group to lend some help here

[fluidnumerics-joe]

@satyanveshd I noticed that with v3.9 of ROCm, this issue still persists with hipcc. Do you have an ETA on when this will be resolved ? Are there any plans to prevent this type of bug from being pushed into future releases ?

[satyanveshd]

The fix didn't make it to 3.9. I believe this should be part of rel 3.10.

[ppanchad-amd]

@fluidnumerics-joe, Sorry for the lack of response. Please try latest ROCm 6.0.2 (HIP 6.0.32831) to see if your issue still exists? If resolved, please close the ticket. Thanks.