Add ROCm support

This PR add support for ROCm.
It is currently based on 7d8ca43532a2c7326d9e74a517cec3fe9eb71fed, I will base it on a more recent commit when cuda_compat is added for atomicAdd of float2, since you will probably need to also implement that for older nvidia GPU as it require compute capability 9.x or higher.
Before being ready for merging, it have two parts that need fixing, they are commented with a "FIXME". First is about compiler flags, pytorch doesn't give us full control of them, from what I tried, gpu-rdc is needed to compile, but pytorch add -fno-gpu-rdc as the last flag. Secondly I had to stub _cuda_raise, I didn't investigate that much about it, but it doesn't work on HIP, a fix or a workaround need to be implemented.

As for performance, on my RX 6700 XT, here is my result on generating 2048 tokens on a 13B model, note that triton based GPTQ is super slow for unknown reason:

GPTQ-for-LLaMa cuda branch: 17.2 tokens/s
GPTQ-for-LLaMa triton branch: 3 tokens/s
exllama: 31.5 tokens/s

I've never had good performance from any of the Triton versions either. Just a bunch of time waiting for the kernels to "warm up" followed by some unsatisfying numbers.
I must admit I have no experience with ROCm, and the last time I had an AMD GPU I can't even remember. I'm a little surprised this is all it takes to get it running. Had to look up what HIP is, and apparently it's just pretty amazing. And I'm not sure how to compare performance across the two architectures, but you're getting half the speed I'm seeing, so either the 6700 XT is very well suited for this, or there's a lot more potential to unlock on the 4090.
The _cuda_raise macro isn't important, there are other ways to check a return value. Don't know why it should cause issues, though, but I'll work around it.
I'm not sure what to do about the compiler flags. I think it's kind of important to avoid any complicated install procedure for the C++ extension, since it seems a lot of people get stuck on that and it might be own my least favorite thing about GPTQ-for-LLaMa, too.
Why is -fgpu-rdc needed, exactly? Just because the code is split across multiple files?

I'm a little surprised this is all it takes to get it running.

Well that because HIP is basically just a clone of CUDA, to make a CUDA project run on ROCm you just have to rename everything to the HIP variant. In our case since it's a cpp extension for pytorch, pytorch automatically run hipify-torch on all .cu files, which rename everything https://github.com/pytorch/pytorch/blob/master/torch/utils/hipify/cuda_to_hip_mappings.py, it doesn't do that for the .h surprisingly. After reading a bit the code, I believe that renaming those .h to .cuh, will allow for hipify-python to also run on them, which would reduce this PR size and increase maintainability. I will try to do that when I have some time.

The _cuda_raise macro isn't important, there are other ways to check a return value. Don't know why it should cause issues, though, but I'll work around it.

I'm not sure, I was a bit fed up because of the issue I had to compile it, so I just stubbed it. It's possible that a simple fix could be made.

Why is -fgpu-rdc needed, exactly? Just because the code is split across multiple files?

I'm not sure, I was rather desperate, with getting lld: error: undefined hidden symbol: __llvm_amdgcn_rcp_f16 while trying to build. I found this similar issue which suggested this fix: https://github.com/ROCm-Developer-Tools/HIP/issues/2196#issuecomment-740165036. AMD doc say that it's to Generates relocatable device code, also known as separate compilation mode, I don't know why the undefined hidden symbol error happen in the first place or why that fix it.
In general, I don't know why pytorch use those compile flags by default, it's the same with __HIP_NO_HALF_CONVERSIONS__, it's always used, thankfully I could disable it, as it's needed for exllama to work. Git blame only point to the PR that implemented HIP compilation for cpp extension, but that didn't have any useful information.

I have rebased locally on latest version. The only thing stopping it being built is HIP missing half2 atomicAdd. I need some help making an efficient implementation that, could you help me with that @turboderp? It require cuda compute capability 6.x or higher so it shouldn't be needed for most nvidia user, but HIP doesn't support it.
In other news, someone shared his benchmark using 2 MI60, getting around 8.8 tokens/s generating 128 tokens on 65B.

I'm willing to test every commit you make as long as it justifies my retarded purchase.

I'm not convinced MI60 is a bad purchase at all. It seems to have a lot of bandwidth, FP16 support... I could see it running very well under the right circumstances.
And I pushed an update for the atomicAdd. I've tested it as much as I can and the function seems to work as it should. To test it you'll need to run with -dq 1 or more, to dequantize at least one layer at load time. Cause there's only one function that does half2 atomicAdd (non-cuBLAS half matmul function) right now, and at the moment it only gets called on pre-dequantized layers.
If it doesn't work with -dq 1, that's not a big loss, though. Turns out it's not a useful feature anyway, and I can do a half version of the matmul to fall back on.

I saw your push by chance, and got real excited. Unfortunately it fails to build.
Regarding the earlier claim of 8+t/s on 65B, it was incorrect. After verifying with test_chatbot.py on ardfork's fork, I only get gibberish, in both single and double GPU mode.
Below is the output from attempting to test inference on your most recent push.
root@a3eef14b80f1:/mnt/textgen/exllama# python test_benchmark_inference.py -d ../models/ausboss_WizardLM-13B-Uncensored-4bit-128g/ -p -ppl -dq 1
Successfully preprocessed all matching files.
Traceback (most recent call last):
  File "/root/miniconda3/lib/python3.10/site-packages/torch/utils/cpp_extension.py", line 1893, in _run_ninja_build
    subprocess.run(
  File "/root/miniconda3/lib/python3.10/subprocess.py", line 526, in run
    raise CalledProcessError(retcode, process.args,
subprocess.CalledProcessError: Command '['ninja', '-v']' returned non-zero exit status 1.

The above exception was the direct cause of the following exception:

Traceback (most recent call last):
  File "/mnt/textgen/exllama/test_benchmark_inference.py", line 1, in <module>
    from model import ExLlama, ExLlamaCache, ExLlamaConfig
  File "/mnt/textgen/exllama/model.py", line 5, in <module>
    import cuda_ext
  File "/mnt/textgen/exllama/cuda_ext.py", line 14, in <module>
    exllama_ext = load(
  File "/root/miniconda3/lib/python3.10/site-packages/torch/utils/cpp_extension.py", line 1284, in load
    return _jit_compile(
  File "/root/miniconda3/lib/python3.10/site-packages/torch/utils/cpp_extension.py", line 1509, in _jit_compile
    _write_ninja_file_and_build_library(
  File "/root/miniconda3/lib/python3.10/site-packages/torch/utils/cpp_extension.py", line 1624, in _write_ninja_file_and_build_library
    _run_ninja_build(
  File "/root/miniconda3/lib/python3.10/site-packages/torch/utils/cpp_extension.py", line 1909, in _run_ninja_build
    raise RuntimeError(message) from e
RuntimeError: Error building extension 'exllama_ext': [1/11] /opt/rocm-5.4.2/bin/hipcc  -DWITH_HIP -DTORCH_EXTENSION_NAME=exllama_ext -DTORCH_API_INCLUDE_EXTENSION_H -DPYBIND11_COMPILER_TYPE=\"_gcc\" -DPYBIND11_STDLIB=\"_libstdcpp\" -DPYBIND11_BUILD_ABI=\"_cxxabi1011\" -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/torch/csrc/api/include -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/TH -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/THC -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/THH -isystem /opt/rocm-5.4.2/include -isystem /opt/rocm-5.4.2/miopen/include -isystem /opt/rocm-5.4.2/hip/include -isystem /root/miniconda3/include/python3.10 -D_GLIBCXX_USE_CXX11_ABI=0 -fPIC -std=c++17 -fPIC -D__HIP_PLATFORM_HCC__=1 -DUSE_ROCM=1 -DCUDA_HAS_FP16=1 -D__HIP_NO_HALF_OPERATORS__=1 -D__HIP_NO_HALF_CONVERSIONS__=1 --amdgpu-target=gfx900 --amdgpu-target=gfx906 --amdgpu-target=gfx908 --amdgpu-target=gfx90a --amdgpu-target=gfx1030 -c /mnt/textgen/exllama/exllama_ext/hip_buffers.hip -o hip_buffers.cuda.o
FAILED: hip_buffers.cuda.o
/opt/rocm-5.4.2/bin/hipcc  -DWITH_HIP -DTORCH_EXTENSION_NAME=exllama_ext -DTORCH_API_INCLUDE_EXTENSION_H -DPYBIND11_COMPILER_TYPE=\"_gcc\" -DPYBIND11_STDLIB=\"_libstdcpp\" -DPYBIND11_BUILD_ABI=\"_cxxabi1011\" -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/torch/csrc/api/include -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/TH -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/THC -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/THH -isystem /opt/rocm-5.4.2/include -isystem /opt/rocm-5.4.2/miopen/include -isystem /opt/rocm-5.4.2/hip/include -isystem /root/miniconda3/include/python3.10 -D_GLIBCXX_USE_CXX11_ABI=0 -fPIC -std=c++17 -fPIC -D__HIP_PLATFORM_HCC__=1 -DUSE_ROCM=1 -DCUDA_HAS_FP16=1 -D__HIP_NO_HALF_OPERATORS__=1 -D__HIP_NO_HALF_CONVERSIONS__=1 --amdgpu-target=gfx900 --amdgpu-target=gfx906 --amdgpu-target=gfx908 --amdgpu-target=gfx90a --amdgpu-target=gfx1030 -c /mnt/textgen/exllama/exllama_ext/hip_buffers.hip -o hip_buffers.cuda.o
Warning: The --amdgpu-target option has been deprecated and will be removed in the future.  Use --offload-arch instead.
Warning: The --amdgpu-target option has been deprecated and will be removed in the future.  Use --offload-arch instead.
Warning: The --amdgpu-target option has been deprecated and will be removed in the future.  Use --offload-arch instead.
Warning: The --amdgpu-target option has been deprecated and will be removed in the future.  Use --offload-arch instead.
Warning: The --amdgpu-target option has been deprecated and will be removed in the future.  Use --offload-arch instead.
In file included from /mnt/textgen/exllama/exllama_ext/hip_buffers.hip:2:
/mnt/textgen/exllama/exllama_ext/cuda_buffers.h:4:10: fatal error: 'cuda_runtime.h' file not found
#include <cuda_runtime.h>
         ^~~~~~~~~~~~~~~~
1 error generated when compiling for gfx1030.
[2/11] /opt/rocm-5.4.2/bin/hipcc  -DWITH_HIP -DTORCH_EXTENSION_NAME=exllama_ext -DTORCH_API_INCLUDE_EXTENSION_H -DPYBIND11_COMPILER_TYPE=\"_gcc\" -DPYBIND11_STDLIB=\"_libstdcpp\" -DPYBIND11_BUILD_ABI=\"_cxxabi1011\" -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/torch/csrc/api/include -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/TH -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/THC -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/THH -isystem /opt/rocm-5.4.2/include -isystem /opt/rocm-5.4.2/miopen/include -isystem /opt/rocm-5.4.2/hip/include -isystem /root/miniconda3/include/python3.10 -D_GLIBCXX_USE_CXX11_ABI=0 -fPIC -std=c++17 -fPIC -D__HIP_PLATFORM_HCC__=1 -DUSE_ROCM=1 -DCUDA_HAS_FP16=1 -D__HIP_NO_HALF_OPERATORS__=1 -D__HIP_NO_HALF_CONVERSIONS__=1 --amdgpu-target=gfx900 --amdgpu-target=gfx906 --amdgpu-target=gfx908 --amdgpu-target=gfx90a --amdgpu-target=gfx1030 -c /mnt/textgen/exllama/exllama_ext/hip_func/half_matmul.hip -o half_matmul.cuda.o
FAILED: half_matmul.cuda.o
/opt/rocm-5.4.2/bin/hipcc  -DWITH_HIP -DTORCH_EXTENSION_NAME=exllama_ext -DTORCH_API_INCLUDE_EXTENSION_H -DPYBIND11_COMPILER_TYPE=\"_gcc\" -DPYBIND11_STDLIB=\"_libstdcpp\" -DPYBIND11_BUILD_ABI=\"_cxxabi1011\" -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/torch/csrc/api/include -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/TH -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/THC -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/THH -isystem /opt/rocm-5.4.2/include -isystem /opt/rocm-5.4.2/miopen/include -isystem /opt/rocm-5.4.2/hip/include -isystem /root/miniconda3/include/python3.10 -D_GLIBCXX_USE_CXX11_ABI=0 -fPIC -std=c++17 -fPIC -D__HIP_PLATFORM_HCC__=1 -DUSE_ROCM=1 -DCUDA_HAS_FP16=1 -D__HIP_NO_HALF_OPERATORS__=1 -D__HIP_NO_HALF_CONVERSIONS__=1 --amdgpu-target=gfx900 --amdgpu-target=gfx906 --amdgpu-target=gfx908 --amdgpu-target=gfx90a --amdgpu-target=gfx1030 -c /mnt/textgen/exllama/exllama_ext/hip_func/half_matmul.hip -o half_matmul.cuda.o
Warning: The --amdgpu-target option has been deprecated and will be removed in the future.  Use --offload-arch instead.
Warning: The --amdgpu-target option has been deprecated and will be removed in the future.  Use --offload-arch instead.
Warning: The --amdgpu-target option has been deprecated and will be removed in the future.  Use --offload-arch instead.
Warning: The --amdgpu-target option has been deprecated and will be removed in the future.  Use --offload-arch instead.
Warning: The --amdgpu-target option has been deprecated and will be removed in the future.  Use --offload-arch instead.
/mnt/textgen/exllama/exllama_ext/hip_func/half_matmul.hip:3:10: fatal error: 'half_matmul.h' file not found
#include "half_matmul.h"
         ^~~~~~~~~~~~~~~
1 error generated when compiling for gfx1030.
[3/11] /opt/rocm-5.4.2/bin/hipcc  -DWITH_HIP -DTORCH_EXTENSION_NAME=exllama_ext -DTORCH_API_INCLUDE_EXTENSION_H -DPYBIND11_COMPILER_TYPE=\"_gcc\" -DPYBIND11_STDLIB=\"_libstdcpp\" -DPYBIND11_BUILD_ABI=\"_cxxabi1011\" -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/torch/csrc/api/include -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/TH -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/THC -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/THH -isystem /opt/rocm-5.4.2/include -isystem /opt/rocm-5.4.2/miopen/include -isystem /opt/rocm-5.4.2/hip/include -isystem /root/miniconda3/include/python3.10 -D_GLIBCXX_USE_CXX11_ABI=0 -fPIC -std=c++17 -fPIC -D__HIP_PLATFORM_HCC__=1 -DUSE_ROCM=1 -DCUDA_HAS_FP16=1 -D__HIP_NO_HALF_OPERATORS__=1 -D__HIP_NO_HALF_CONVERSIONS__=1 --amdgpu-target=gfx900 --amdgpu-target=gfx906 --amdgpu-target=gfx908 --amdgpu-target=gfx90a --amdgpu-target=gfx1030 -c /mnt/textgen/exllama/exllama_ext/hip_func/rope.hip -o rope.cuda.o
FAILED: rope.cuda.o
/opt/rocm-5.4.2/bin/hipcc  -DWITH_HIP -DTORCH_EXTENSION_NAME=exllama_ext -DTORCH_API_INCLUDE_EXTENSION_H -DPYBIND11_COMPILER_TYPE=\"_gcc\" -DPYBIND11_STDLIB=\"_libstdcpp\" -DPYBIND11_BUILD_ABI=\"_cxxabi1011\" -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/torch/csrc/api/include -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/TH -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/THC -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/THH -isystem /opt/rocm-5.4.2/include -isystem /opt/rocm-5.4.2/miopen/include -isystem /opt/rocm-5.4.2/hip/include -isystem /root/miniconda3/include/python3.10 -D_GLIBCXX_USE_CXX11_ABI=0 -fPIC -std=c++17 -fPIC -D__HIP_PLATFORM_HCC__=1 -DUSE_ROCM=1 -DCUDA_HAS_FP16=1 -D__HIP_NO_HALF_OPERATORS__=1 -D__HIP_NO_HALF_CONVERSIONS__=1 --amdgpu-target=gfx900 --amdgpu-target=gfx906 --amdgpu-target=gfx908 --amdgpu-target=gfx90a --amdgpu-target=gfx1030 -c /mnt/textgen/exllama/exllama_ext/hip_func/rope.hip -o rope.cuda.o
Warning: The --amdgpu-target option has been deprecated and will be removed in the future.  Use --offload-arch instead.
Warning: The --amdgpu-target option has been deprecated and will be removed in the future.  Use --offload-arch instead.
Warning: The --amdgpu-target option has been deprecated and will be removed in the future.  Use --offload-arch instead.
Warning: The --amdgpu-target option has been deprecated and will be removed in the future.  Use --offload-arch instead.
Warning: The --amdgpu-target option has been deprecated and will be removed in the future.  Use --offload-arch instead.
/mnt/textgen/exllama/exllama_ext/hip_func/rope.hip:3:10: fatal error: 'rope.h' file not found
#include "rope.h"
         ^~~~~~~~
1 error generated when compiling for gfx1030.
[4/11] /opt/rocm-5.4.2/bin/hipcc  -DWITH_HIP -DTORCH_EXTENSION_NAME=exllama_ext -DTORCH_API_INCLUDE_EXTENSION_H -DPYBIND11_COMPILER_TYPE=\"_gcc\" -DPYBIND11_STDLIB=\"_libstdcpp\" -DPYBIND11_BUILD_ABI=\"_cxxabi1011\" -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/torch/csrc/api/include -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/TH -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/THC -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/THH -isystem /opt/rocm-5.4.2/include -isystem /opt/rocm-5.4.2/miopen/include -isystem /opt/rocm-5.4.2/hip/include -isystem /root/miniconda3/include/python3.10 -D_GLIBCXX_USE_CXX11_ABI=0 -fPIC -std=c++17 -fPIC -D__HIP_PLATFORM_HCC__=1 -DUSE_ROCM=1 -DCUDA_HAS_FP16=1 -D__HIP_NO_HALF_OPERATORS__=1 -D__HIP_NO_HALF_CONVERSIONS__=1 --amdgpu-target=gfx900 --amdgpu-target=gfx906 --amdgpu-target=gfx908 --amdgpu-target=gfx90a --amdgpu-target=gfx1030 -c /mnt/textgen/exllama/exllama_ext/hip_func/q4v2_recons.hip -o q4v2_recons.cuda.o
FAILED: q4v2_recons.cuda.o
/opt/rocm-5.4.2/bin/hipcc  -DWITH_HIP -DTORCH_EXTENSION_NAME=exllama_ext -DTORCH_API_INCLUDE_EXTENSION_H -DPYBIND11_COMPILER_TYPE=\"_gcc\" -DPYBIND11_STDLIB=\"_libstdcpp\" -DPYBIND11_BUILD_ABI=\"_cxxabi1011\" -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/torch/csrc/api/include -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/TH -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/THC -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/THH -isystem /opt/rocm-5.4.2/include -isystem /opt/rocm-5.4.2/miopen/include -isystem /opt/rocm-5.4.2/hip/include -isystem /root/miniconda3/include/python3.10 -D_GLIBCXX_USE_CXX11_ABI=0 -fPIC -std=c++17 -fPIC -D__HIP_PLATFORM_HCC__=1 -DUSE_ROCM=1 -DCUDA_HAS_FP16=1 -D__HIP_NO_HALF_OPERATORS__=1 -D__HIP_NO_HALF_CONVERSIONS__=1 --amdgpu-target=gfx900 --amdgpu-target=gfx906 --amdgpu-target=gfx908 --amdgpu-target=gfx90a --amdgpu-target=gfx1030 -c /mnt/textgen/exllama/exllama_ext/hip_func/q4v2_recons.hip -o q4v2_recons.cuda.o
Warning: The --amdgpu-target option has been deprecated and will be removed in the future.  Use --offload-arch instead.
Warning: The --amdgpu-target option has been deprecated and will be removed in the future.  Use --offload-arch instead.
Warning: The --amdgpu-target option has been deprecated and will be removed in the future.  Use --offload-arch instead.
Warning: The --amdgpu-target option has been deprecated and will be removed in the future.  Use --offload-arch instead.
Warning: The --amdgpu-target option has been deprecated and will be removed in the future.  Use --offload-arch instead.
/mnt/textgen/exllama/exllama_ext/hip_func/q4v2_recons.hip:3:10: fatal error: 'q4v2_recons.h' file not found
#include "q4v2_recons.h"
         ^~~~~~~~~~~~~~~
1 error generated when compiling for gfx1030.
[5/11] /opt/rocm-5.4.2/bin/hipcc  -DWITH_HIP -DTORCH_EXTENSION_NAME=exllama_ext -DTORCH_API_INCLUDE_EXTENSION_H -DPYBIND11_COMPILER_TYPE=\"_gcc\" -DPYBIND11_STDLIB=\"_libstdcpp\" -DPYBIND11_BUILD_ABI=\"_cxxabi1011\" -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/torch/csrc/api/include -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/TH -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/THC -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/THH -isystem /opt/rocm-5.4.2/include -isystem /opt/rocm-5.4.2/miopen/include -isystem /opt/rocm-5.4.2/hip/include -isystem /root/miniconda3/include/python3.10 -D_GLIBCXX_USE_CXX11_ABI=0 -fPIC -std=c++17 -fPIC -D__HIP_PLATFORM_HCC__=1 -DUSE_ROCM=1 -DCUDA_HAS_FP16=1 -D__HIP_NO_HALF_OPERATORS__=1 -D__HIP_NO_HALF_CONVERSIONS__=1 --amdgpu-target=gfx900 --amdgpu-target=gfx906 --amdgpu-target=gfx908 --amdgpu-target=gfx90a --amdgpu-target=gfx1030 -c /mnt/textgen/exllama/exllama_ext/hip_func/q4v2_sequential.hip -o q4v2_sequential.cuda.o
FAILED: q4v2_sequential.cuda.o
/opt/rocm-5.4.2/bin/hipcc  -DWITH_HIP -DTORCH_EXTENSION_NAME=exllama_ext -DTORCH_API_INCLUDE_EXTENSION_H -DPYBIND11_COMPILER_TYPE=\"_gcc\" -DPYBIND11_STDLIB=\"_libstdcpp\" -DPYBIND11_BUILD_ABI=\"_cxxabi1011\" -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/torch/csrc/api/include -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/TH -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/THC -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/THH -isystem /opt/rocm-5.4.2/include -isystem /opt/rocm-5.4.2/miopen/include -isystem /opt/rocm-5.4.2/hip/include -isystem /root/miniconda3/include/python3.10 -D_GLIBCXX_USE_CXX11_ABI=0 -fPIC -std=c++17 -fPIC -D__HIP_PLATFORM_HCC__=1 -DUSE_ROCM=1 -DCUDA_HAS_FP16=1 -D__HIP_NO_HALF_OPERATORS__=1 -D__HIP_NO_HALF_CONVERSIONS__=1 --amdgpu-target=gfx900 --amdgpu-target=gfx906 --amdgpu-target=gfx908 --amdgpu-target=gfx90a --amdgpu-target=gfx1030 -c /mnt/textgen/exllama/exllama_ext/hip_func/q4v2_sequential.hip -o q4v2_sequential.cuda.o
Warning: The --amdgpu-target option has been deprecated and will be removed in the future.  Use --offload-arch instead.
Warning: The --amdgpu-target option has been deprecated and will be removed in the future.  Use --offload-arch instead.
Warning: The --amdgpu-target option has been deprecated and will be removed in the future.  Use --offload-arch instead.
Warning: The --amdgpu-target option has been deprecated and will be removed in the future.  Use --offload-arch instead.
Warning: The --amdgpu-target option has been deprecated and will be removed in the future.  Use --offload-arch instead.
/mnt/textgen/exllama/exllama_ext/hip_func/q4v2_sequential.hip:3:10: fatal error: 'q4v2_sequential.h' file not found
#include "q4v2_sequential.h"
         ^~~~~~~~~~~~~~~~~~~
1 error generated when compiling for gfx1030.
[6/11] c++ -MMD -MF exllama_ext_hip.o.d -DTORCH_EXTENSION_NAME=exllama_ext -DTORCH_API_INCLUDE_EXTENSION_H -DPYBIND11_COMPILER_TYPE=\"_gcc\" -DPYBIND11_STDLIB=\"_libstdcpp\" -DPYBIND11_BUILD_ABI=\"_cxxabi1011\" -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/torch/csrc/api/include -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/TH -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/THC -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/THH -isystem /opt/rocm-5.4.2/include -isystem /opt/rocm-5.4.2/miopen/include -isystem /opt/rocm-5.4.2/hip/include -isystem /root/miniconda3/include/python3.10 -D_GLIBCXX_USE_CXX11_ABI=0 -fPIC -std=c++17 -c /mnt/textgen/exllama/exllama_ext/exllama_ext_hip.cpp -o exllama_ext_hip.o -fPIC -D__HIP_PLATFORM_HCC__=1 -DUSE_ROCM=1
FAILED: exllama_ext_hip.o
c++ -MMD -MF exllama_ext_hip.o.d -DTORCH_EXTENSION_NAME=exllama_ext -DTORCH_API_INCLUDE_EXTENSION_H -DPYBIND11_COMPILER_TYPE=\"_gcc\" -DPYBIND11_STDLIB=\"_libstdcpp\" -DPYBIND11_BUILD_ABI=\"_cxxabi1011\" -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/torch/csrc/api/include -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/TH -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/THC -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/THH -isystem /opt/rocm-5.4.2/include -isystem /opt/rocm-5.4.2/miopen/include -isystem /opt/rocm-5.4.2/hip/include -isystem /root/miniconda3/include/python3.10 -D_GLIBCXX_USE_CXX11_ABI=0 -fPIC -std=c++17 -c /mnt/textgen/exllama/exllama_ext/exllama_ext_hip.cpp -o exllama_ext_hip.o -fPIC -D__HIP_PLATFORM_HCC__=1 -DUSE_ROCM=1
In file included from /mnt/textgen/exllama/exllama_ext/exllama_ext_hip.cpp:10:
/mnt/textgen/exllama/exllama_ext/cuda_buffers.h:4:10: fatal error: cuda_runtime.h: No such file or directory
    4 | #include <cuda_runtime.h>
      |          ^~~~~~~~~~~~~~~~
compilation terminated.
[7/11] /opt/rocm-5.4.2/bin/hipcc  -DWITH_HIP -DTORCH_EXTENSION_NAME=exllama_ext -DTORCH_API_INCLUDE_EXTENSION_H -DPYBIND11_COMPILER_TYPE=\"_gcc\" -DPYBIND11_STDLIB=\"_libstdcpp\" -DPYBIND11_BUILD_ABI=\"_cxxabi1011\" -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/torch/csrc/api/include -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/TH -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/THC -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/THH -isystem /opt/rocm-5.4.2/include -isystem /opt/rocm-5.4.2/miopen/include -isystem /opt/rocm-5.4.2/hip/include -isystem /root/miniconda3/include/python3.10 -D_GLIBCXX_USE_CXX11_ABI=0 -fPIC -std=c++17 -fPIC -D__HIP_PLATFORM_HCC__=1 -DUSE_ROCM=1 -DCUDA_HAS_FP16=1 -D__HIP_NO_HALF_OPERATORS__=1 -D__HIP_NO_HALF_CONVERSIONS__=1 --amdgpu-target=gfx900 --amdgpu-target=gfx906 --amdgpu-target=gfx908 --amdgpu-target=gfx90a --amdgpu-target=gfx1030 -c /mnt/textgen/exllama/exllama_ext/hip_func/column_remap.hip -o column_remap.cuda.o
FAILED: column_remap.cuda.o
/opt/rocm-5.4.2/bin/hipcc  -DWITH_HIP -DTORCH_EXTENSION_NAME=exllama_ext -DTORCH_API_INCLUDE_EXTENSION_H -DPYBIND11_COMPILER_TYPE=\"_gcc\" -DPYBIND11_STDLIB=\"_libstdcpp\" -DPYBIND11_BUILD_ABI=\"_cxxabi1011\" -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/torch/csrc/api/include -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/TH -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/THC -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/THH -isystem /opt/rocm-5.4.2/include -isystem /opt/rocm-5.4.2/miopen/include -isystem /opt/rocm-5.4.2/hip/include -isystem /root/miniconda3/include/python3.10 -D_GLIBCXX_USE_CXX11_ABI=0 -fPIC -std=c++17 -fPIC -D__HIP_PLATFORM_HCC__=1 -DUSE_ROCM=1 -DCUDA_HAS_FP16=1 -D__HIP_NO_HALF_OPERATORS__=1 -D__HIP_NO_HALF_CONVERSIONS__=1 --amdgpu-target=gfx900 --amdgpu-target=gfx906 --amdgpu-target=gfx908 --amdgpu-target=gfx90a --amdgpu-target=gfx1030 -c /mnt/textgen/exllama/exllama_ext/hip_func/column_remap.hip -o column_remap.cuda.o
Warning: The --amdgpu-target option has been deprecated and will be removed in the future.  Use --offload-arch instead.
Warning: The --amdgpu-target option has been deprecated and will be removed in the future.  Use --offload-arch instead.
Warning: The --amdgpu-target option has been deprecated and will be removed in the future.  Use --offload-arch instead.
Warning: The --amdgpu-target option has been deprecated and will be removed in the future.  Use --offload-arch instead.
Warning: The --amdgpu-target option has been deprecated and will be removed in the future.  Use --offload-arch instead.
/mnt/textgen/exllama/exllama_ext/hip_func/column_remap.hip:3:10: fatal error: 'column_remap.h' file not found
#include "column_remap.h"
         ^~~~~~~~~~~~~~~~
1 error generated when compiling for gfx1030.
[8/11] /opt/rocm-5.4.2/bin/hipcc  -DWITH_HIP -DTORCH_EXTENSION_NAME=exllama_ext -DTORCH_API_INCLUDE_EXTENSION_H -DPYBIND11_COMPILER_TYPE=\"_gcc\" -DPYBIND11_STDLIB=\"_libstdcpp\" -DPYBIND11_BUILD_ABI=\"_cxxabi1011\" -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/torch/csrc/api/include -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/TH -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/THC -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/THH -isystem /opt/rocm-5.4.2/include -isystem /opt/rocm-5.4.2/miopen/include -isystem /opt/rocm-5.4.2/hip/include -isystem /root/miniconda3/include/python3.10 -D_GLIBCXX_USE_CXX11_ABI=0 -fPIC -std=c++17 -fPIC -D__HIP_PLATFORM_HCC__=1 -DUSE_ROCM=1 -DCUDA_HAS_FP16=1 -D__HIP_NO_HALF_OPERATORS__=1 -D__HIP_NO_HALF_CONVERSIONS__=1 --amdgpu-target=gfx900 --amdgpu-target=gfx906 --amdgpu-target=gfx908 --amdgpu-target=gfx90a --amdgpu-target=gfx1030 -c /mnt/textgen/exllama/exllama_ext/hip_func/rms_norm.hip -o rms_norm.cuda.o
FAILED: rms_norm.cuda.o
/opt/rocm-5.4.2/bin/hipcc  -DWITH_HIP -DTORCH_EXTENSION_NAME=exllama_ext -DTORCH_API_INCLUDE_EXTENSION_H -DPYBIND11_COMPILER_TYPE=\"_gcc\" -DPYBIND11_STDLIB=\"_libstdcpp\" -DPYBIND11_BUILD_ABI=\"_cxxabi1011\" -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/torch/csrc/api/include -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/TH -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/THC -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/THH -isystem /opt/rocm-5.4.2/include -isystem /opt/rocm-5.4.2/miopen/include -isystem /opt/rocm-5.4.2/hip/include -isystem /root/miniconda3/include/python3.10 -D_GLIBCXX_USE_CXX11_ABI=0 -fPIC -std=c++17 -fPIC -D__HIP_PLATFORM_HCC__=1 -DUSE_ROCM=1 -DCUDA_HAS_FP16=1 -D__HIP_NO_HALF_OPERATORS__=1 -D__HIP_NO_HALF_CONVERSIONS__=1 --amdgpu-target=gfx900 --amdgpu-target=gfx906 --amdgpu-target=gfx908 --amdgpu-target=gfx90a --amdgpu-target=gfx1030 -c /mnt/textgen/exllama/exllama_ext/hip_func/rms_norm.hip -o rms_norm.cuda.o
Warning: The --amdgpu-target option has been deprecated and will be removed in the future.  Use --offload-arch instead.
Warning: The --amdgpu-target option has been deprecated and will be removed in the future.  Use --offload-arch instead.
Warning: The --amdgpu-target option has been deprecated and will be removed in the future.  Use --offload-arch instead.
Warning: The --amdgpu-target option has been deprecated and will be removed in the future.  Use --offload-arch instead.
Warning: The --amdgpu-target option has been deprecated and will be removed in the future.  Use --offload-arch instead.
/mnt/textgen/exllama/exllama_ext/hip_func/rms_norm.hip:3:10: fatal error: 'rms_norm.h' file not found
#include "rms_norm.h"
         ^~~~~~~~~~~~
1 error generated when compiling for gfx1030.
[9/11] /opt/rocm-5.4.2/bin/hipcc  -DWITH_HIP -DTORCH_EXTENSION_NAME=exllama_ext -DTORCH_API_INCLUDE_EXTENSION_H -DPYBIND11_COMPILER_TYPE=\"_gcc\" -DPYBIND11_STDLIB=\"_libstdcpp\" -DPYBIND11_BUILD_ABI=\"_cxxabi1011\" -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/torch/csrc/api/include -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/TH -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/THC -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/THH -isystem /opt/rocm-5.4.2/include -isystem /opt/rocm-5.4.2/miopen/include -isystem /opt/rocm-5.4.2/hip/include -isystem /root/miniconda3/include/python3.10 -D_GLIBCXX_USE_CXX11_ABI=0 -fPIC -std=c++17 -fPIC -D__HIP_PLATFORM_HCC__=1 -DUSE_ROCM=1 -DCUDA_HAS_FP16=1 -D__HIP_NO_HALF_OPERATORS__=1 -D__HIP_NO_HALF_CONVERSIONS__=1 --amdgpu-target=gfx900 --amdgpu-target=gfx906 --amdgpu-target=gfx908 --amdgpu-target=gfx90a --amdgpu-target=gfx1030 -c /mnt/textgen/exllama/exllama_ext/hip_func/q4v2_matmul.hip -o q4v2_matmul.cuda.o
FAILED: q4v2_matmul.cuda.o
/opt/rocm-5.4.2/bin/hipcc  -DWITH_HIP -DTORCH_EXTENSION_NAME=exllama_ext -DTORCH_API_INCLUDE_EXTENSION_H -DPYBIND11_COMPILER_TYPE=\"_gcc\" -DPYBIND11_STDLIB=\"_libstdcpp\" -DPYBIND11_BUILD_ABI=\"_cxxabi1011\" -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/torch/csrc/api/include -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/TH -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/THC -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/THH -isystem /opt/rocm-5.4.2/include -isystem /opt/rocm-5.4.2/miopen/include -isystem /opt/rocm-5.4.2/hip/include -isystem /root/miniconda3/include/python3.10 -D_GLIBCXX_USE_CXX11_ABI=0 -fPIC -std=c++17 -fPIC -D__HIP_PLATFORM_HCC__=1 -DUSE_ROCM=1 -DCUDA_HAS_FP16=1 -D__HIP_NO_HALF_OPERATORS__=1 -D__HIP_NO_HALF_CONVERSIONS__=1 --amdgpu-target=gfx900 --amdgpu-target=gfx906 --amdgpu-target=gfx908 --amdgpu-target=gfx90a --amdgpu-target=gfx1030 -c /mnt/textgen/exllama/exllama_ext/hip_func/q4v2_matmul.hip -o q4v2_matmul.cuda.o
Warning: The --amdgpu-target option has been deprecated and will be removed in the future.  Use --offload-arch instead.
Warning: The --amdgpu-target option has been deprecated and will be removed in the future.  Use --offload-arch instead.
Warning: The --amdgpu-target option has been deprecated and will be removed in the future.  Use --offload-arch instead.
Warning: The --amdgpu-target option has been deprecated and will be removed in the future.  Use --offload-arch instead.
Warning: The --amdgpu-target option has been deprecated and will be removed in the future.  Use --offload-arch instead.
/mnt/textgen/exllama/exllama_ext/hip_func/q4v2_matmul.hip:3:10: fatal error: 'q4v2_matmul.h' file not found
#include "q4v2_matmul.h"
         ^~~~~~~~~~~~~~~
1 error generated when compiling for gfx1030.
[10/11] /opt/rocm-5.4.2/bin/hipcc  -DWITH_HIP -DTORCH_EXTENSION_NAME=exllama_ext -DTORCH_API_INCLUDE_EXTENSION_H -DPYBIND11_COMPILER_TYPE=\"_gcc\" -DPYBIND11_STDLIB=\"_libstdcpp\" -DPYBIND11_BUILD_ABI=\"_cxxabi1011\" -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/torch/csrc/api/include -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/TH -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/THC -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/THH -isystem /opt/rocm-5.4.2/include -isystem /opt/rocm-5.4.2/miopen/include -isystem /opt/rocm-5.4.2/hip/include -isystem /root/miniconda3/include/python3.10 -D_GLIBCXX_USE_CXX11_ABI=0 -fPIC -std=c++17 -fPIC -D__HIP_PLATFORM_HCC__=1 -DUSE_ROCM=1 -DCUDA_HAS_FP16=1 -D__HIP_NO_HALF_OPERATORS__=1 -D__HIP_NO_HALF_CONVERSIONS__=1 --amdgpu-target=gfx900 --amdgpu-target=gfx906 --amdgpu-target=gfx908 --amdgpu-target=gfx90a --amdgpu-target=gfx1030 -c /mnt/textgen/exllama/exllama_ext/hip_func/q4v2_mlp.hip -o q4v2_mlp.cuda.o
FAILED: q4v2_mlp.cuda.o
/opt/rocm-5.4.2/bin/hipcc  -DWITH_HIP -DTORCH_EXTENSION_NAME=exllama_ext -DTORCH_API_INCLUDE_EXTENSION_H -DPYBIND11_COMPILER_TYPE=\"_gcc\" -DPYBIND11_STDLIB=\"_libstdcpp\" -DPYBIND11_BUILD_ABI=\"_cxxabi1011\" -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/torch/csrc/api/include -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/TH -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/THC -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/THH -isystem /opt/rocm-5.4.2/include -isystem /opt/rocm-5.4.2/miopen/include -isystem /opt/rocm-5.4.2/hip/include -isystem /root/miniconda3/include/python3.10 -D_GLIBCXX_USE_CXX11_ABI=0 -fPIC -std=c++17 -fPIC -D__HIP_PLATFORM_HCC__=1 -DUSE_ROCM=1 -DCUDA_HAS_FP16=1 -D__HIP_NO_HALF_OPERATORS__=1 -D__HIP_NO_HALF_CONVERSIONS__=1 --amdgpu-target=gfx900 --amdgpu-target=gfx906 --amdgpu-target=gfx908 --amdgpu-target=gfx90a --amdgpu-target=gfx1030 -c /mnt/textgen/exllama/exllama_ext/hip_func/q4v2_mlp.hip -o q4v2_mlp.cuda.o
Warning: The --amdgpu-target option has been deprecated and will be removed in the future.  Use --offload-arch instead.
Warning: The --amdgpu-target option has been deprecated and will be removed in the future.  Use --offload-arch instead.
Warning: The --amdgpu-target option has been deprecated and will be removed in the future.  Use --offload-arch instead.
Warning: The --amdgpu-target option has been deprecated and will be removed in the future.  Use --offload-arch instead.
Warning: The --amdgpu-target option has been deprecated and will be removed in the future.  Use --offload-arch instead.
/mnt/textgen/exllama/exllama_ext/hip_func/q4v2_mlp.hip:3:10: fatal error: 'q4v2_mlp.h' file not found
#include "q4v2_mlp.h"
         ^~~~~~~~~~~~
1 error generated when compiling for gfx1030.
ninja: build stopped: subcommand failed.

(base) root@a3eef14b80f1:/mnt/textgen/exllama#

Also, if you want me to attempt merging and rebasing you have to explain it to me like I am a retard.
I have never worked with git more than git clone and git pull. :)

Hey, I'm not very good at git either.
But, from the output it looks like it's due to that thing where PyTorch doesn't want to HIPify .h tiles, so I tried renaming them all to .cuh instead. The C++ compiler doesn't care anyway.
Still, I have no way to test this myself, so who knows what will happen.

so I tried renaming them all to .cuh instead.

Attempted again, but received identical errors below (truncated to last build since they are all identical):
[11/12] /opt/rocm-5.4.2/bin/hipcc  -DWITH_HIP -DTORCH_EXTENSION_NAME=exllama_ext -DTORCH_API_INCLUDE_EXTENSION_H -DPYBIND11_COMPILER_TYPE=\"_gcc\" -DPYBIND11_STDLIB=\"_libstdcpp\" -DPYBIND11_BUILD_ABI=\"_cxxabi1011\" -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/torch/csrc/api/include -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/TH -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/THC -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/THH -isystem /opt/rocm-5.4.2/include -isystem /opt/rocm-5.4.2/miopen/include -isystem /opt/rocm-5.4.2/hip/include -isystem /root/miniconda3/include/python3.10 -D_GLIBCXX_USE_CXX11_ABI=0 -fPIC -std=c++17 -fPIC -D__HIP_PLATFORM_HCC__=1 -DUSE_ROCM=1 -DCUDA_HAS_FP16=1 -D__HIP_NO_HALF_OPERATORS__=1 -D__HIP_NO_HALF_CONVERSIONS__=1 --amdgpu-target=gfx900 --amdgpu-target=gfx906 --amdgpu-target=gfx908 --amdgpu-target=gfx90a --amdgpu-target=gfx1030 -c /mnt/textgen/exllama/exllama_ext/hip_func/column_remap.hip -o column_remap.cuda.o
FAILED: column_remap.cuda.o
/opt/rocm-5.4.2/bin/hipcc  -DWITH_HIP -DTORCH_EXTENSION_NAME=exllama_ext -DTORCH_API_INCLUDE_EXTENSION_H -DPYBIND11_COMPILER_TYPE=\"_gcc\" -DPYBIND11_STDLIB=\"_libstdcpp\" -DPYBIND11_BUILD_ABI=\"_cxxabi1011\" -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/torch/csrc/api/include -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/TH -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/THC -isystem /root/miniconda3/lib/python3.10/site-packages/torch/include/THH -isystem /opt/rocm-5.4.2/include -isystem /opt/rocm-5.4.2/miopen/include -isystem /opt/rocm-5.4.2/hip/include -isystem /root/miniconda3/include/python3.10 -D_GLIBCXX_USE_CXX11_ABI=0 -fPIC -std=c++17 -fPIC -D__HIP_PLATFORM_HCC__=1 -DUSE_ROCM=1 -DCUDA_HAS_FP16=1 -D__HIP_NO_HALF_OPERATORS__=1 -D__HIP_NO_HALF_CONVERSIONS__=1 --amdgpu-target=gfx900 --amdgpu-target=gfx906 --amdgpu-target=gfx908 --amdgpu-target=gfx90a --amdgpu-target=gfx1030 -c /mnt/textgen/exllama/exllama_ext/hip_func/column_remap.hip -o column_remap.cuda.o
Warning: The --amdgpu-target option has been deprecated and will be removed in the future.  Use --offload-arch instead.
Warning: The --amdgpu-target option has been deprecated and will be removed in the future.  Use --offload-arch instead.
Warning: The --amdgpu-target option has been deprecated and will be removed in the future.  Use --offload-arch instead.
Warning: The --amdgpu-target option has been deprecated and will be removed in the future.  Use --offload-arch instead.
Warning: The --amdgpu-target option has been deprecated and will be removed in the future.  Use --offload-arch instead.
/mnt/textgen/exllama/exllama_ext/hip_func/column_remap.hip:3:10: fatal error: 'column_remap.cuh' file not found
#include "column_remap.cuh"
         ^~~~~~~~~~~~~~~~~~
1 error generated when compiling for gfx1030.
ninja: build stopped: subcommand failed.

(base) root@a3eef14b80f1:/mnt/textgen/exllama#

Looks like it's HIPifying them, but the #include "column_remap.cuh" line is calling to the un-HIPified version.
Am I reading too far into this?

after checking a little bit to see what went well in ardfork's version, I found that the exllama_ext.cpp was being HIPified a little strangely, and I think it's due to the location change to using a new folder for cuda_fuct.
Below is the first 15 lines from exllama_ext.cpp from ardfork's repo:
#include <torch/extension.h>
#include <c10/cuda/CUDAGuard.h>
#include <cuda_runtime.h>
#include <cuda_fp16.h>
#include <cstdint>
#include <cstdio>

#include "column_remap.h"
#include "q4v2_matmul.h"
#include "q4v2_mlp.h"
#include "q4v2_recons.h"
#include "q4v2_sequential.h"
#include "rms_norm.h"
#include "util.h"

They have all of these calls for files within the same directory.
For the up-to-date branch on your repo,  they are imported as so:
#include <torch/extension.h>
#include <c10/cuda/CUDAGuard.h>
#include <ATen/cuda/CUDAContext.h>
#include <cuda_runtime.h>
#include <cuda_fp16.h>
#include <cstdint>
#include <cstdio>

#include "cpu_func/rep_penalty.h"

#include "cuda_buffers.cuh"
#include "cuda_func/column_remap.cuh"
#include "cuda_func/half_matmul.cuh"
#include "cuda_func/q4v2_matmul.cuh"
#include "cuda_func/q4v2_mlp.cuh"
#include "cuda_func/q4v2_recons.cuh"
#include "cuda_func/q4v2_sequential.cuh"
#include "cuda_func/rms_norm.cuh"
#include "cuda_func/rope.cuh"
#include "util.cuh"

That's cool, but HIPify creates a new dir called hip_func, and places these HIPified functions in there, but still attempts to call them from cuda_func:
#include <torch/extension.h>
#include <ATen/hip/impl/HIPGuardImplMasqueradingAsCUDA.h>
#include <ATen/hip/HIPContext.h>
#include <hip/hip_runtime.h>
#include <hip/hip_fp16.h>
#include <cstdint>
#include <cstdio>

#include "cpu_func/rep_penalty.h"

#include "cuda_buffers.cuh"
#include "cuda_func/column_remap.cuh"
#include "cuda_func/half_matmul.cuh"
#include "cuda_func/q4v2_matmul.cuh"
#include "cuda_func/q4v2_mlp.cuh"
#include "cuda_func/q4v2_recons.cuh"
#include "cuda_func/q4v2_sequential.cuh"
#include "cuda_func/rms_norm.cuh"
#include "cuda_func/rope.cuh"
#include "util.cuh"```

Is this an error in HIPify, or is it an issue with file pathing and C/Python fuckery?

Pulled a lot of stuff out of ardfork's rework, involving a lot of manual renaming and the classic copypasta technique. Nearly there but running into only math errors.
I hope someone smarter than me can do math.
In file included from /mnt/textgen/exllama/exllama_ext/hip_func/rope.hip:5:
/mnt/textgen/exllama/exllama_ext/hip_func/../matrix.cuh:218:10: error: no viable conversion from 'unsigned int' to 'half' (aka '__half')
    half result1_ = __hadd(result1.x, result1.y);

You could try removing the dot_product_8_dual and dot_product_8_dual_buffered as they're not referenced anywhere right now. I can't say why it's producing an error about unsigned ints when adding two fp16 values and storing to another fp16. But those two functions at least you can comment out and disregard.

I pushed my rebase.
But it doesn't work, it compile but during inferring, it SIGABRT with "Cannot find Symbol", setting AMD_LOG_LEVEL to try to debug show:
:3:hip_module.cpp           :469 :<PHONE_NUMBER> us: 12874: [tid:0x7f00dce85740]  hipLaunchKernel ( 0x7efe2bfb1a20, {80,480,1}, {16,4,1}, 0x7ffe178f3210, 0, stream:<null> )
:3:devprogram.cpp           :2676:<PHONE_NUMBER> us: 12874: [tid:0x7f00dce85740] Using Code Object V4.
:3:devprogram.cpp           :2979:<PHONE_NUMBER> us: 12874: [tid:0x7f00dce85740] For Init/Fini: Kernel Name: _Z52ncclKernelLL128Debug_SendRecv_RING_SIMPLE_Sum_int8_tP11ncclDevCommmP8ncclWork
:3:devprogram.cpp           :2979:<PHONE_NUMBER> us: 12874: [tid:0x7f00dce85740] For Init/Fini: Kernel Name: _Z47ncclKernelLL128_SendRecv_RING_SIMPLE_Sum_int8_tP11ncclDevCommmP8ncclWork
:3:devprogram.cpp           :2979:<PHONE_NUMBER> us: 12874: [tid:0x7f00dce85740] For Init/Fini: Kernel Name: _Z47ncclKernelDebug_SendRecv_RING_SIMPLE_Sum_int8_tP11ncclDevCommmP8ncclWork
:3:devprogram.cpp           :2979:<PHONE_NUMBER> us: 12874: [tid:0x7f00dce85740] For Init/Fini: Kernel Name: _Z42ncclKernel_SendRecv_RING_SIMPLE_Sum_int8_tP11ncclDevCommmP8ncclWork
:1:hip_global.cpp           :97  :<PHONE_NUMBER> us: 12874: [tid:0x7f00dce85740] Cannot find Symbol with name: _Z27rms_norm_row_product_kernelP6__halfPfii

I don't really understand why it fail while launching kernel, if we remove this call, it will fail on another kernel (silu_mul_cuda_kernel). Those lines were here before, I don't see any significant change that would break things. Other similar project with similar codebase work well when hipified, like GPTQ for example, and it also launch kernel the same way. I verified and exllama_ext.so correctly have those symbols, I'm afraid I don't know what is the cause of this and I don't have contact with any person knowledgeable with HIP.

It may be the -fno-gpu-rdc flag. It could be that the HIPifying code just doesn't know how to deal with linking multiple files. You could try creating a single .cu file that just #includes all the others, and then just have that one file (along with the .cpp files) at the top of cuda_ext.py.

"Cannot find Symbol"

Same here. Sorry I couldn't help more.

I'm considering pulling the trigger and returning my two MI60s before I'm stuck with them for life... Thinking about buying A4000s to fill the 4 PCIe slots the MI60s hold currently, and would bring me to the same total VRAM.
Am I being too pussy, or is ROCm just doomed?

It may be the -fno-gpu-rdc flag.

This seems to to be the correct assessment. Trying to build even a trivial extension without this flag causes SIGABRT "Cannot find Symbol": seems like it was there for a good reason.
I'm not entirely sure if adding this flag initially solved the problem either: I suspect it's likely that it simply moved compile-time linking errors into runtime. Then, these errors manifested in _cuda_raise which was stubbed. By stubbing the macro, I think it may have been removing all kernel calls, resulting in no errors and unusually high speeds.
The error itself (undefined hidden symbol: __llvm_amdgcn_rcp_f16) is strange, but seems to only be caused the call to h2rcp(sum) in q4v2_mlp.cu. I tried a naïve approach of calculating the reciprocal by hand:
half2 r;
r.x = 1.0 / sum.x;
r.y = 1.0 / sum.y

This compiles, but chat mode still seems to give gibberish for me. I'm not sure if that's because of this or something else not working.

This compiles, but chat mode still seems to give gibberish for me.

Ardfork's original pull req had this issue too. I was benchmarking good speeds, but had not thought to check perplexity nor chatbot.
Chatbot just dumps trash and perplexity fails.

Oddly enough, the inference benchmark seems to mostly work for me now. It shows a perplexity of 5.6797 which is about in line with the 13B model I used for it. I'm not sure if this is some coincidence or because of some difference between the two, such as config.

Could you include your changes in a PR or maybe just a line num + edit? I'd like to give it a try, for continuity.

The only changes made was undoing any changes to pytorch regarding -fno-gpu-rdc and removing -fgpu-rdc from the extra_cuda_cflags list in cuda_exy.py. You can probably just change -fgpu-rdc to -fno-gpu-rdc in there though and skip fixing torch.
The only other things is the reciprocal change to q4v2_mlp.cu I noted above:
@@ -29 +29,3 @@
-    half2 r = h2rcp(sum);
+    half2 r;
+    r.x = 1.0 / sum.x;
+    r.y = 1.0 / sum.y;

Which I'm not sure is correct, and is probably slower than ideal even if it is.

I'm surprised that compiles. But if there's a problem with the fused MLP you could try running with -mlp normal, then it shouldn't use any code in q4v2_mlp.cu. Also you can try `-mm pytorch_only' which should avoid using q4v2_matmul.cu. It won't run great like that, but there would be less CUDA code.

I'm surprised that compiles.

Is it that bad 😄
I'll freely admit this is not my usual domain.

you could try running with -mlp normal

This seems to have no effect on things, perplexity doesn't change and chatbot outputs are similar.

Also you can try -mm pytorch_only

This doesn't seem to affect perplexity, but in chat mode it makes things a little more "sane?" Still gibberish, but a lot less of repeating the same tokens and shorter outputs. For reference, it does indeed run at about 2/3 of the usual speed.

We got too cocky...

Are we so back???

I can confirm that setting both at once provides what seems like correct output. Since this disables what sounds like a lot of the main optimizations of exllama, the low speeds are not unexpected. I would speculate that the earlier speed of about 3t/s is a more reasonable estimate of performance if the rest can be made to work.

I can confirm that setting both at once provides what seems like correct output. Since this disables what sounds like a lot of the main optimizations of exllama, the low speeds are not unexpected. I would speculate that the earlier speed of about 3t/s is a more reasonable estimate of performance if the rest can be made to work.

Actually, I seem to get much higher speeds natively on the ooba webui.

If it runs at all with -mm pytorch_only that means it's at least able to run some HIPified CUDA code, which is a start. It does seem like it's the half2 stuff that's causing issues. Which makes sense if the AMD GPUs don't have the same SIMD features as NVIDIA and half2 is being emulated. The fact that some half2 intrinsics like h2exp compile while others like h2rcp don't also seems to support this.
I'm currently putting the finishing touches on an initial version of the web UI (I ended up making it a bit more advanced than I originally intended), and after that I'm going to look at some more optimization and building out the C++ extension a lot. I need multiple versions of every CUDA function, at least one tuned for 40-series GPUs and one for 30-series as it turns out the two have somewhat different profiles. Then there will be a more involved setup with possibly an autotuning option.
And with that, making sure each function also has a version that avoids half2 types will be trivial. It also creates a place to plug in dedicated AMD-optimized alternatives, although I'd have a hard time writing those without the hardware to test them.

The error itself (undefined hidden symbol: __llvm_amdgcn_rcp_f16) is strange, but seems to only be caused the call to h2rcp(sum) in q4v2_mlp.cu. I tried a naïve approach of calculating the reciprocal by hand:
half2 r;
r.x = 1.0 / sum.x;
r.y = 1.0 / sum.y

I investigated a bit further, and it's actually already been fixed upstream, it should be fixed in ROCm 5.5.0 but my distro does not yet package it. I will see if I can backport the fix.

This compiles, but chat mode still seems to give gibberish for me. I'm not sure if that's because of this or something else not working.

Sadly, I'm also not sure what that happen, others hipified project like llama.cpp, bitsandbytes, gptq all work correctly without any hacks, I wonder what is going wrong here.

should be fixed in ROCm 5.5.0

I'm on ROCm 5.5.0, actually. (though, it's merely in an ubuntu docker)

https://github.com/ROCm-Developer-Tools/hipamd/commit/36f929bb42a252533a61f1e275de6565bfbaa7ae
It should have been part of 5.4.4. I'm on 5.4.3, so it isn't fixed on my machine. And most people are probably using 5.4.2 as that what the official pytorch ROCm use.

Pushed a fix for h2rcp. Now we only have to find where the gibberish output come from.

On my machine, -mm pytorch_only make it work,

Is this with or without the earlier edit to cpp_extension.py?
Mine was reverted, and I get gibberish no matter the -mm pytorch_only or not.

Is this with or without the earlier edit to cpp_extension.py?

It is with a clean venv, clone of my rocm branch with no edits.
My only test was launching with python test_chatbot.py -d $model -un "Maxine" -p prompt_assistant.txt -nnl -temp 1.0 -topp .75 -mm pytorch_only and typing "What is the capital of Russia?\n". It answered coherently with "Chatbort: The capital of Russia is Moscow.".
But now that I test again with a different model, it doesn't work. The first one where it worked was something quantized with GPTQ triton branch, with Groupsize (inferred): 128 and Act-order (inferred): yes. The second one, which didn't work was quantized with cuda branch, with Groupsize (inferred): 128 and Act-order (inferred): no.
@turboderp: do you have any idea what kind of functions is used in common in such case, one being in matmul and the other in model quantized with GPTQ cuda branch?

Act-order models should run strictly more CUDA code than the other kind. So it's very weird if those are the ones that work. If you could run with the -dbg flag, maybe? Dump the output on pastebin, and I'll have a look.

If you could run with the -dbg flag, maybe? Dump the output on pastebin, and I'll have a look.
Here a dump of python test_benchmark_inference.py -d $model -dbg, with -mm pytorch_only and without on both models: https://gist.github.com/ardfork/5995e627ae70a6c2d927a2b097804752

The only one that work in chatbot is the case "gptq-triton-mm_pytorch_only".

Okay. So what sets the working case apart from all the others is that the act-order matrices would override switching to the CUDA MLP, and -mm pytorch_only disables the quantized matmul elsewhere. So the CUDA functions that seem to work are:

column_remap_cuda
q4v2_recons_cuda
q4v2_sequential_cuda
rms_norm_cuda
rope_cuda

Definitely broken with HIP:

q4v2_matmul_cuda
q4v2_mlp_cuda (possibly because it calls q4v2_matmul_cuda)

And maybe working (can't tell from these examples):

half_matmul_cuda

Now, many of the other CUDA implementations allow for  Torch datatypes (c10::Half etc.) in the CUDA code, which may be why it's better able to HIPify half2 types and intrinsics. Whereas I cast everything in the C++ part to avoid having to include the Torch headers in the CUDA parts. Tomorrow I think I'm going to rewrite at least the matmul code to work on half types, optionally, and then we'll see from there.

It's only in develop branch, so maybe for 5.6.

Have you tried to compile hip from the develop branch? Didn't quite work too well for me, but I'm stuck inside with docker and my IQ is low, so I can't confirm.

Now, many of the other CUDA implementations allow for Torch datatypes (c10::Half etc.) in the CUDA code, which may be why it's better able to HIPify half2 types and intrinsics. Whereas I cast everything in the C++ part to avoid having to include the Torch headers in the CUDA parts. Tomorrow I think I'm going to rewrite at least the matmul code to work on half types, optionally, and then we'll see from there.

The thing is that pytorch itself use the same script to make it work with ROCm, they run hipify-python. They of course have some #if defined(USE_ROCM) to fix some stuff, but if something doesn't work, it could be interesting to see if pytorch also use it and check how they make it work.

I've been investigating the possibility of hipifying the header files automatically for a bit, and it seems that .cuh or .h doesn't matter. There is a way to make it work however: setting build_directory = "." in the call to load(…) seems to pull in the headers, hipify them, and not include them as sources (which was something I tried previously, breaking builds).
This has the upside of removing the #ifdef clutter at the cost of increased build artifact clutter. Hipifying already caused a fair bit of clutter, so it's not so bad from that perspective, but I imagine those building with cuda might not like the sudden presence of build artifacts.
If others can confirm this works, then it could be a good way to reduce maintenance until more rocm-specific code shows up, and reduce the diff of this PR somewhat substantially.

Not sure where y'all are at with this right now, but in case it helps, I've been adding tuning arguments today. Lots more to come, but currently starting with -nh2 should switch all the kernels to versions that don't rely on half2 types. So potentially no more issues with h2rcp and whatnot. Not sure if it helps, but it's something at least. There should also be a couple of places where you could remove the half2 code altogether with #ifdefs, if it maybe helps with compilation.

currently starting with -nh2 should switch all the kernels to versions that don't rely on half2 types. Not sure if it helps, but it's something at least.

That was it! It currently work on all my models with -nh2. Thank you.
The speed is a bit slower than my initial test where it was spouting gibberish, but it is still faster than GPTQ (and use less VRAM). Here are the result on my machine, as a reminder it's with an RX 6700 XT (and an i7-7700K):
 ** Time, Inference: 3.70 seconds
 ** Speed: 518.37 tokens/second
 -- Generating 128 tokens, 1920 token prompt...
 ** Speed: 23.25 tokens/second
 -- Generating 128 tokens, 4 token prompt...
 ** Speed: 28.99 tokens/second

Would be interesting to check if the problem is on all half2 stuff or only in specific one, it works with only --rmsnorm_no_half2 --matmul_no_half2 --silu_no_half2. That mean that --rope_no_half2 doesn't change anything. Does it mean that half2 rope work or simply that it doesn't use rope in my case?

So potentially no more issues with h2rcp and whatnot.

There was the same problem with hrcp blocking compilation that I fixed in the same way as h2rcp.

There should also be a couple of places where you could remove the half2 code altogether with #ifdefs, if it maybe helps with compilation.

I don't think I want to remove all the half2 code all together, maybe it will work on a more recent version of ROCm or on a more recent GPU architecture. What I would prefer is to use no half2 by default (in HIP build) and still be able to force half2 with a different arg.

I've been investigating the possibility of hipifying the header files automatically for a bit, and it seems that .cuh or .h doesn't matter. There is a way to make it work however: setting build_directory = "." in the call to load(…) seems to pull in the headers, hipify them, and not include them as sources (which was something I tried previously, breaking builds).

I also have tried multiple way at first to make them hipify automatically (without calling hipify-python myself before pytorch cpp_extension) but didn't manage to make it work. I tried your way using build_directory = ".", but couldn't make it work either, because of some unicode error if I remember correctly. Maybe it's because I have my venv in my project dir. Could you send a patch or include more information on how you made it work?

I happened to be working on updating things at the same time and can confirm all of your observations on what makes it work :+1:.

I also have tried multiple way at first to make them hipify automatically (without calling hipify-python myself before pytorch cpp_extension) but didn't manage to make it work. I tried your way using build_directory = ".", but couldn't make it work either, because of some unicode error if I remember correctly. Maybe it's because I have my venv in my project dir. Could you send a patch or include more information on how you made it work?

I did at some point make some changes to pytorch, but I'm pretty sure I reverted those. I don't think that there was anything still lingering there, but it might be possible. I do think I saw that it payed attention to git configuration while I was looking into it, so it's possible that I wasn't encountering the error because I have a .gitignore file locally:
*hip*
*.o
*.so
Pipfile*
*ninja*
.venv
__pycache__
!exllama_ext/hip_compat.cuh

If torch does use this to ignore the venv, then that could explain why it worked for me. I did try removing .venv from that list and it still seemed to work, however. I'm not sure if yours is actually just venv and it being a hidden file changes things, or if our installed packages differ in some way. The specific error you get may be useful.

it's possible that I wasn't encountering the error because I have a .gitignore file locally

I also have a local .gitignore, but it only has my venv in it to avoid some tools greping my venv when I just want to search in the project.

The specific error you get may be useful.

I tried again just editing cuda_ext.py on master and get UnicodeDecodeError: 'utf-8' codec can't decode byte 0x89 in position 0: invalid start byte, which from a quick internet search look like the start of a PNG file. A few PNG are in my venv from some python packages (cmake, networkx, sympy, werkzeug).
Looking at cpp_extension and hipify_python source I'm not sure if that the right play, I guess it work because cpp_extension call hipify with project_directory=build_dir,. What I don't understand is that they have some test for cpp_extension that also have some .cuh headers that contain some cuda import so they should fail if hipify doesn't convert them. The only difference is that in their test, they add the dir containing .cu and .cuh to extra_include_paths, which is then given to hipify as header_include_dirs=(extra_include_paths if extra_include_paths is not None else []), and from a quick look at hipify_python source, it should try to find these headers and hipify them. But simply adding exllama_ext in those header include dirs doesn't  seem to work.

Guys, could someone point me which resources to look our if I'd like to grasp some meaningful CUDA and ROCm knowledge needed to program AI math? Assuming one has some Python / C++ knowledge but don't know much about GPU programming, shaders, tensor cores and similar black magic.

So is it more or less in a working state now, just to summarize?

I think an appropriate summary is that it works with --rmsnorm_no_half2 --matmul_no_half2 --silu_no_half2 and that it may be possible to remove if defined(USE_ROCM) from most header files, but the configuration required to make this work well eludes us still.

But simply adding exllama_ext in those header include dirs doesn't seem to work.

I examined the code a bit, and it seems like the pytorch code may be a bit buggy: it tries to absolute the paths relative to the output directory rather than the project directory (which is what I would have expected). Fortunately, you can provide the absolute path as with the sources:
extra_include_paths = [os.path.join(library_dir, "exllama_ext")]

This works for me. Again though, others should verify for the same reasons as before.

extra_include_paths = [os.path.join(library_dir, "exllama_ext")]

This works for me. Again though, others should verify for the same reasons as before.

Seem to also work on my case, I will rework my patch to use that method, it's cleaner and should be easier to maintain. I also want to check a bit about my change with cublasHgemm, ideally it should use hipblasHgemm and not use rocblas_hgemm directly, AMD doc and hipify-perl use it, so it should work. I'm guessing that python doesn't use it and hipify_python is not really an external/supported project but only used with pytorch as far as I know.
If everything use HIP, it should theoretically also work with NVIDIA GPU, that might seem useless, but that useful if you have multiple GPU composed of AMD and NVIDIA.
After that, it should be ready to merge after testing that it doesn't break CUDA build. Automatically disabling half2 could be useful, but it's not needed for this PR to be useful.

Very happy to report that I'm managing to run a 33B model using two AMD GPUs in a 16GB+8GB configuration. Speeds are very nice too, well in excess of what I was getting with GPU offloading in llama.cpp/similar.
GPU 1: Radeon RX 6800 XT
GPU 2: Radeon RX 5700 XT
VRAM Split: 12,7
Model: llama-30b-supercot-4bit-128g-cuda
 -- Inference, first pass.
 ** Time, Inference: 24.26 seconds
 ** Speed: 79.14 tokens/second
 -- Generating 128 tokens, 1920 token prompt...
 ** Speed: 8.93 tokens/second
 -- Generating 128 tokens, 4 token prompt...
 ** Speed: 13.39 tokens/second
 -- Loading dataset...
 -- Testing..........
 ** Perplexity: 4.7271

I did have to go through the process of building pytorch manually to support the 5700 XT (gfx 1010), but other than that it seems to work well.

Hmm, I was writing a message saying that the PR was finished but now it crash with Memory access fault by GPU node-1 (Agent handle: 0xxxxxxxxxxxxx) on address 0xxxxxxxxxxxxx. Reason: Page not present or supervisor privilege.
The only thing I did was rebase on master, everything was working correctly before, seem like there was a regression somewhere.

Memory access fault by GPU node-1 (Agent handle: 0xxxxxxxxxxxxx) on address 0xxxxxxxxxxxxx. Reason: Page not present or supervisor privilege.

I've actually gotten this error a ton with my MI60s, no amount of research has turned up a concrete fix.
I haven't had it in a while, maybe turn it off and back on again?

 It was the fault of 167f60193f9e9eb8c31d2cbd26ff3ddebe25b8bd. Please @turboderp don't hide code change in an "Update todo" commit.

Alright, I quickly fixed that, don't know why it was causing issue.
Now that PR should be good. It is ready to be reviewed and merged.

Alright, I quickly fixed that, don't know why it was causing issue.
Now that PR should be good. It is ready to be reviewed and merged.

(base) root@a3eef14b80f1:/mnt/textgen/exllama# python test_benchmark_inference.py -d ../models/Neko-Institute-of-Science_LLaMA-65B-4bit-128g/ -p -ppl -gs 20,20
Successfully preprocessed all matching files.
 -- Tokenizer: ../models/Neko-Institute-of-Science_LLaMA-65B-4bit-128g/tokenizer.model
 -- Model config: ../models/Neko-Institute-of-Science_LLaMA-65B-4bit-128g/config.json
 -- Model: ../models/Neko-Institute-of-Science_LLaMA-65B-4bit-128g/llama-65b-4bit-128g.safetensors
 -- Sequence length: 2048
 -- Tuning:
 -- --matmul_recons_thd: 8
 -- --fused_mlp_thd: 2
 -- --sdp_thd: 8
 -- --rmsnorm_no_half2
 -- --rope_no_half2
 -- --matmul_no_half2
 -- --silu_no_half2
 -- Options: ['gpu_split: 20,20', 'perf', 'perplexity']
 ** Time, Load model: 48.31 seconds
 ** Time, Load tokenizer: 0.02 seconds
 -- Groupsize (inferred): 128
 -- Act-order (inferred): yes
 ** VRAM, Model: [cuda:0] 20,901.84 MB - [cuda:1] 12,563.61 MB
 -- Warmup pass 1...
 ** Time, Warmup: 24.68 seconds
 -- Warmup pass 2...
 ** Time, Warmup: 24.23 seconds
 -- Warmup pass 3...
 ** Time, Warmup: 24.24 seconds
 -- Inference, first pass.
 ** Time, Inference: 24.32 seconds
 ** Speed: 78.94 tokens/second
 -- Generating 128 tokens, 1920 token prompt...
 ** Speed: 4.10 tokens/second
 -- Generating 128 tokens, 4 token prompt...
 ** Speed: 4.81 tokens/second
 ** VRAM, Inference: [cuda:0] 4,284.22 MB - [cuda:1] 2,890.08 MB
 ** VRAM, Total: [cuda:0] 25,186.06 MB - [cuda:1] 15,453.68 MB
 -- Loading dataset...
 -- Testing..........
 ** Perplexity: 4.1895
(base) root@a3eef14b80f1:/mnt/textgen/exllama#

Looks good! Working with your new commit, here's 65B on two MI60.

Yeah sorry about that. PyCharm makes it a little too easy to commit and push changes. I'll be more careful.

Well, this is great. Thanks guys. I'm still a little amazed it works at all with so few changes at the end of the day.
Based on some of these benchmarks I suspect there are lots of AMD specific optimizations to look at, too. The prompt speed on the MI60 is really low, for instance, compared to the speed per token which is kind of decent for a GPU from 2018. Seems the kernels are 4-5 times slower than my 4090+3090 setup, while rocBLAS is about 13 times slower. Seems off, but who knows.

By the way, the next big optimization I'm going for will be using CUDA Graphs. And "as of my knowledge cutoff in September 2021" there's no ROCm equivalent. Does anyone know if this has changed since? If not I'll try to keep the regular CUDA stuff as is and implement graphs within a #ifndef USE_ROCM scope from the start so we won't have to keep coming back to it.

It might work, was added in around May 2021. And was released in ROCm 5.3 which was released in October 2022, most people are probably using >= 5.4.2.
In pytorch they use a bunch of #if !defined(USE_ROCM) || ROCM_VERSION >= 50300

I've started work on moving to graphs. It's in the "graphs" branch for now, since I'd like someone to at least try it out on ROCm before I start breaking stuff for people.

I've started work on moving to graphs. It's in the "graphs" branch for now, since I'd like someone to at least try it out on ROCm before I start breaking stuff for people.

Some are not handled by hipify_python, I made a small patch, do you want me to open a PR for it?
diff --git a/exllama_ext/cuda_func/rms_norm.cu b/exllama_ext/cuda_func/rms_norm.cu
index d327a65..34f643e 100644
--- a/exllama_ext/cuda_func/rms_norm.cu
+++ b/exllama_ext/cuda_func/rms_norm.cu
@@ -3,6 +3,17 @@
 #include "../util.cuh"
 #include "../matrix.cuh"

+#if defined(USE_ROCM)
+#define cudaGraphAddKernelNode hipGraphAddKernelNode
+#define cudaGraphAddMemsetNode hipGraphAddMemsetNode
+#define cudaGraphCreate hipGraphCreate
+#define cudaGraphExecKernelNodeSetParams hipGraphExecKernelNodeSetParams
+#define cudaGraphExecMemsetNodeSetParams hipGraphExecMemsetNodeSetParams
+#define cudaGraphNode_t hipGraphNode_t
+#define cudaKernelNodeParams hipKernelNodeParams
+#define cudaMemsetParams hipMemsetParams
+#endif
+
 const int THREADS_X = 32;
 const int THREADS_Y = 8;
 const int BLOCKSIZE_X = 16;
@@ -207,8 +218,18 @@ void rms_norm_cuda

     void* args1[] = { &x, &temp, &rows, &dim };
     void* args2[] = { &x, &w, &out, &temp, &epsilon, &r_dim, &rows, &dim };
-    cudaKernelNodeParams params1 = { (void*) rms_norm_kernel_func1, blocks, threads, 0, args1, nullptr };
-    cudaKernelNodeParams params2 = { (void*) rms_norm_kernel_func2, blocks, threads, 0, args2, nullptr };
+    cudaKernelNodeParams params1 = { .func           = (void *)rms_norm_kernel_func1,
+                                     .gridDim        = blocks,
+                                     .blockDim       = threads,
+                                     .sharedMemBytes = 0,
+                                     .kernelParams   = args1,
+                                     .extra          = nullptr };
+    cudaKernelNodeParams params2 = { .func           = (void *)rms_norm_kernel_func2,
+                                     .gridDim        = blocks,
+                                     .blockDim       = threads,
+                                     .sharedMemBytes = 0,
+                                     .kernelParams   = args2,
+                                     .extra          = nullptr };

     if (!contexts[device_index].rms_norm_graph_init)
     {

It is working correctly with a small test.  But seem a bit slower:
 ** Time, Inference: 3.71 seconds
 ** Speed: 516.85 tokens/second
 -- Generating 128 tokens, 1920 token prompt...
 ** Speed: 22.12 tokens/second
 -- Generating 128 tokens, 4 token prompt...
 ** Speed: 26.88 tokens/second

Also, I want to detail a bit my patch making process.
First I try to run it, if they are compilation errors, I note the files which have some. In our case, it's rms_norm.hip.
Secondly I run hipify-perl -examine $file on the hipified (by pytorch) file, hipify-perl is a bit more generic instead of being a tool specific for pytorch. In our case, it results in this:
  warning: exllama/exllama_ext/hip_func/rms_norm.hip:79: unsupported device function "__shfl_down_sync": //         acc += __shfl_down_sync(0xffffffff, acc, offset);

[HIPIFY] info: file 'exllama/exllama_ext/hip_func/rms_norm.hip' statistics:
  CONVERTED refs count: 31
  TOTAL lines of code: 270
  WARNINGS: 1
[HIPIFY] info: CONVERTED refs by names:
  cudaGraphAddKernelNode => hipGraphAddKernelNode: 3
  cudaGraphAddMemsetNode => hipGraphAddMemsetNode: 2
  cudaGraphCreate => hipGraphCreate: 2
  cudaGraphExecKernelNodeSetParams => hipGraphExecKernelNodeSetParams: 3
  cudaGraphExecMemsetNodeSetParams => hipGraphExecMemsetNodeSetParams: 2
  cudaGraphNode_t => hipGraphNode_t: 4
  cudaKernelNodeParams => hipKernelNodeParams: 3
  cudaMemsetParams => hipMemsetParams: 2

I then copy those cuda => hip line into the .cu files, put them inside #if defined(USE_ROCM) and use a simple sed s/  \(.\+\) => \(.\+\): [0-9]\+/#define \1 \2.
Now, I try again to run it, and fix the remaining errors. In our case, that was with difference between cudaMemsetParams and hipMemsetParams. They have the same members in their struct but not in the same order. We can easily fix that by using designated initializer style.

Thanks, I added the changes.
As for the speed, I'm hoping it's not because HIP emulates graphs too clumsily, e.g. by destroying and recreating them whenever parameters change, which would defeat the purpose somewhat. But maybe it's just constant overhead and it will start to flip with longer graphs.

Well the nice thing is that HIP is open source so you can actually check. The less nice thing is that not many people have time to read such complex code.

These are the only errors I'm still getting, didn't get a chance to test ardfork's own changes and ended up just pulling most recent commit.
In file included from /mnt/textgen/exllama/exllama_ext/hip_func/rms_norm.hip:8:
/mnt/textgen/exllama/exllama_ext/hip_func/../hip_compat.cuh:19:26: error: unknown type name 'hipblasStatus_t'
__host__ __forceinline__ hipblasStatus_t __compat_hipblasHgemm(hipblasHandle_t    handle,
                         ^
/mnt/textgen/exllama/exllama_ext/hip_func/../hip_compat.cuh:19:64: error: unknown type name 'hipblasHandle_t'
__host__ __forceinline__ hipblasStatus_t __compat_hipblasHgemm(hipblasHandle_t    handle,
                                                               ^
/mnt/textgen/exllama/exllama_ext/hip_func/../hip_compat.cuh:20:64: error: unknown type name 'hipblasOperation_t'
                                                               hipblasOperation_t transA,
                                                               ^
/mnt/textgen/exllama/exllama_ext/hip_func/../hip_compat.cuh:21:64: error: unknown type name 'hipblasOperation_t'
                                                               hipblasOperation_t transB,
                                                               ^
/mnt/textgen/exllama/exllama_ext/hip_func/../hip_compat.cuh:34:48: error: unknown type name 'hipblasHalf'
                        reinterpret_cast<const hipblasHalf *>(alpha),
                                               ^
/mnt/textgen/exllama/exllama_ext/hip_func/../hip_compat.cuh:35:48: error: unknown type name 'hipblasHalf'
                        reinterpret_cast<const hipblasHalf *>(AP), lda,
                                               ^
/mnt/textgen/exllama/exllama_ext/hip_func/../hip_compat.cuh:36:48: error: unknown type name 'hipblasHalf'
                        reinterpret_cast<const hipblasHalf *>(BP), ldb,
                                               ^
/mnt/textgen/exllama/exllama_ext/hip_func/../hip_compat.cuh:37:48: error: unknown type name 'hipblasHalf'
                        reinterpret_cast<const hipblasHalf *>(beta),
                                               ^
/mnt/textgen/exllama/exllama_ext/hip_func/../hip_compat.cuh:38:42: error: unknown type name 'hipblasHalf'
                        reinterpret_cast<hipblasHalf *>(CP), ldc);
                                         ^

Yes, he changed my patch a bit, which cause this error.
Also, compiling with -O3 give me a ~2.5% speed boost.

Damn it. I'll put it in the .cu file for now, but I really would prefer to have the translation fixes in one place rather than spread throughout the CUDA code, because there's a lot of refactoring coming. Is there a way I can force it to generate HIP code without an AMD GPU just to check if it still compiles?

Is there a way I can force it to generate HIP code without an AMD GPU just to check if it still compiles?

Actually HIP is supposed to run on both NVIDIA and AMD. And it's supposed to do so without any performance loss when running NVIDIA, as it's just a wrapper for CUDA function. I don't know if you need the whole ROCm stack, probably not. Try just installing pytorch ROCm in a new venv: pip install torch torchvision torchaudio --extra-index-url https://download.pytorch.org/whl/rocm5.4.2. If that isn't enough, either install ROCm stuff with your distro or try the docker route.

but I really would prefer to have the translation fixes in one place rather than spread throughout the CUDA code

Well in that case, you would just need to add #include <hipblas/hipblas.h> in hip_compat.cuh I guess.

Here's 65B again, on new "graph" branch:
-- Inference, first pass.
 ** Time, Inference: 24.43 seconds
 ** Speed: 78.59 tokens/second
 -- Generating 128 tokens, 1920 token prompt...
 ** Speed: 3.99 tokens/second
 -- Generating 128 tokens, 4 token prompt...
 ** Speed: 4.64 tokens/second
 ** VRAM, Inference: [cuda:0] 4,284.22 MB - [cuda:1] 2,890.08 MB
 ** VRAM, Total: [cuda:0] 25,185.82 MB - [cuda:1] 15,453.44 MB
 -- Loading dataset...
 -- Testing..........
 ** Perplexity: 4.1894

at least it's consistent, right?

Hi @ardfork , I don't get why half2 is disabled by default when compiling for rocm. Could you explain?

Hi @ardfork , I don't get why half2 is disabled by default when compiling for rocm. Could you explain?

Because it is broken. Search "gibberish" in this thread for more information. Also check out #146 for a fix.

Thank you, missed this one! I hit the same issue.

In this PR for me cuda_compat.cuh gets renamed by the hipifier to hip_compat.cuh which conflicts with the file of the same name. I guess this is not intended, is it?
Given this in cuda_compat.cuh:
#if defined(__CUDA_ARCH__) || defined(USE_ROCM)
#if __CUDA_ARCH__ < 700 || defined(USE_ROCM)

__device__ __forceinline__ void atomicAdd(half* address, half val) { atomicAdd_half(address, val); }

#if __CUDA_ARCH__ < 600 || defined(USE_ROCM)
__device__ __forceinline__ void atomicAdd(half2* address, half2 val) { atomicAdd_half2(address, val); }
#endif

#endif
#endif

In this PR for me cuda_compat.cuh gets renamed by the hipifier to hip_compat.cuh which conflicts with the file of the same name. I guess this is not intended, is it?
Given this in cuda_compat.cuh:
#if defined(__CUDA_ARCH__) || defined(USE_ROCM)
#if __CUDA_ARCH__ < 700 || defined(USE_ROCM)

__device__ __forceinline__ void atomicAdd(half* address, half val) { atomicAdd_half(address, val); }

#if __CUDA_ARCH__ < 600 || defined(USE_ROCM)
__device__ __forceinline__ void atomicAdd(half2* address, half2 val) { atomicAdd_half2(address, val); }
#endif

#endif
#endif

Without renaming this file I am getting a no matching function for call to 'atomicAdd' error.

atomicAdd for half isn't supported on RDNA2/3 by the way.


MONDO:0013710 colorectal cancer, hereditary nonpolyposis, type 5

colorectal cancer, hereditary nonpolyposis, type 5 (MONDO:0013710)
should have the parent
MONDO:0005575 | colorectal cancer
(or maybe the  parentage should be via Lynch syndrome which is already a parent of MONDO:0013710)?

From NCIT 06/11/20:
I am almost certain Lynch Syndrome should NOT be a subclass of colorectal cancer but let me confirm with our Pathologist.

So I'm a bit confused now.
If hereditary nonpolyposis colorectal cancer-5
and the hereditary nonpolyposis colorectal cancers can't go under colorectal cancers,
(and these are the most common form of hereditary colon cancers)
which diseases can go under colorectal cancer ?
Or can "hereditary nonpolyposis colorectal cancer-5"
have the cancer parent  because it is used to describe the colon  cancer,
but the syndrome itself can't have the parent because it can  include other types of cancer?
I'm curious because  the current situation gives inconsistencies like this:
familial colorectal cancer (MONDO:0023113)
for example is defined

Definition
Parents

Familial colon cancer is a cluster of colon cancer within a family. Most cases ofcolon cancer occur sporadically in people with little to no family history of the condition. Approximately 3-5% of colon cancer is considered 'hereditary' and is thought to be caused by an inherited predisposition tocolon cancer that is passed down through a family in an autosomal dominant or autosomal recessive manner. In some of these families, the underlying genetic cause is not known; however, many of these cases are caused by changes ( mutations ) in the APC , MYH , MLH1 , MSH2 , MSH6 , PMS2 , EPCAM , PTEN , STK11 , SMAD4 , BMPR1A , NTHL1 , POLE , and POLD1 genes (which are associated with hereditary cancer syndromes). An additional 10-30% of people diagnosed with colon cancer have a significant family history of the condition but have no identifiable mutation in a gene known to cause a hereditary predisposition to colon cancer. These clusters of colon cancer are likely due to a combination of gene(s) and other shared factors such as environment and lifestyle. High-risk cancer screening and other preventative measures such as prophylactic surgeries are typically recommended in people who have an increased risk for colon cancer based on their personal and/or family histories.

which I presume means that it  would be correct  to annotate MSH6
the causative gene for
colorectal cancer, hereditary nonpolyposis, type 5 (MONDO:0013710)
to colorectal cancer (hence the parentage suggestion)

From the NCIT pathologist:
This syndrome manifests primarily with hereditary nonpolyposis colon cancer due to germline mutations of the DNA mismatch repair genes. But, importantly patients with this syndrome develop endometrial, ovarian, renal pelvis, gastric, and small intestinal adenocarcinomas as well. I would urge them to place it under the term Syndrome.

Currently 'Lynch syndrome' is a child of 'hereditary nonpolyposis colon cancer', which is defined as:
Hereditary nonpolyposis colon cancer (HNPCC) is a cancer-predisposing condition characterized by the development of colorectal cancer not associated with colorectal polyposis, endometrial cancer, and various other cancers (such as malignant epithelial tumor of ovary, gastric, biliary tract, small bowel, and urinary tract cancer) that are frequently diagnosed at an early age.
So this parent seems incorrect.
We are thinking about obsoleting 'syndromic disease' though (see https://github.com/monarch-initiative/mondo/issues/1642)
I think Lynch syndrome should be a child of 'hereditary neoplastic syndrome', based on the definition.

@ValWood

Or can "hereditary nonpolyposis colorectal cancer-5"
have the cancer parent because it is used to describe the colon cancer,
but the syndrome itself can't have the parent because it can include other types of cancer?

I just asked the NCIt pathologist about placing Lynch syndrome as a child of colorectal cancer (and cc'd you on the email.)
colorectal cancer, hereditary nonpolyposis, type 5 (MONDO:0013710) is not in NCIt, so I'm not sure if I can ask them about it.
In OMIM, this term is defined as a cancer predisposition syndrome characterized by onset of colorectal cancer and/or extracolonic cancers, particularly endometrial cancer, usually in mid-adulthood.
So I'm not sure MONDO:0005575 colorectal cancer is the correct parent.
I'm not totally sure. @cmungall could you weigh in?

Everywhere I looked Lynch syndrome and Hereditary nonpolyposis colon cancer (HNPCC)
are synonymous?
https://en.wikipedia.org/wiki/Hereditary_nonpolyposis_colorectal_cancer#:~:text=Hereditary nonpolyposis colorectal cancer (HNPCC,%2C%20upper%20urinary%20tract%2C%20brain%2C
https://ghr.nlm.nih.gov/condition/lynch-syndrome
https://www.ouh.nhs.uk/patient-guide/leaflets/files/10055Plynchsyndrome.pdf

https://my.clevelandclinic.org/health/diseases/17195-lynch-syndrome--hnpcc
describes the difference
Is Lynch syndrome the same as hereditary non-polyposis colorectal cancer (HNPCC)?
Sometimes Lynch syndrome is also referred to as hereditary non-polyposis colorectal cancer (HNPCC), but the two conditions are subtly different.
Lynch syndrome is classified by a mutation in the MMR gene. HNPCC is defined by family history and individuals must meet specific criteria, which are:
At least three relatives have colorectal or other HNPCC-related cancers, with one person being a first-degree relative (parent, child, or sibling) of the other two;
At least two generations have been affected;
At least one person received a diagnosis of colorectal cancer before age 50;
Familial adenomatous polyposis has been excluded.
So the logic issue  that  prompted my ticket is  that currently the
colorectal cancer, hereditary nonpolyposis, type 5
etc are under "lynch syndrome"
They need to be housed instead under the HNPCC term.
Comments would be useful to describe the difference because the causative proteins will require separate annotations to both Lynch syndrome and   the appropriate HNPCC term. The HNPCC terms should presumably somehow have the colorectal parent.
somehow have MONDO:0005575 | colorectal cancer parent)

Hi @ValWood

colorectal cancer, hereditary nonpolyposis, type 5
etc are under "lynch syndrome"
They need to be housed instead under the HNPCC term.

I can remove "Lynch syndrome" as a superclass from 'colorectal cancer, hereditary nonpolyposis, type 5'. ('hereditary nonpolyposis colon cancer' is already a parent for this class).
What do you mean by "etc"?
Should types 2, 4, 6, 7 and 8 also be children of 'hereditary nonpolyposis colon cancer' and not 'Lynch syndrome'?

Comments would be useful to describe the difference because the causative proteins will require separate annotations to both Lynch syndrome and the appropriate HNPCC term. The HNPCC terms should presumably somehow have the colorectal parent.
somehow have MONDO:0005575 | colorectal cancer parent)

Can you be more specific - what comment would you like us to add and to which class? I'm happy to do so.
Are you asking to add MONDO:0005575 colorectal cancer as a parent off 'hereditary nonpolyposis colon cancer'
Thanks for clarifying!

Should types 2, 4, 6, 7 and 8 also be children of 'hereditary nonpolyposis colon cancer' and not 'Lynch syndrome'?

yes
"Lynch syndrome"  needs a comment
"Lynch syndrome is classified by a mutation in the MMR gene. " (Or adding to the definition)
It seems that describing as
HNPCC is quite complicated
"
defined by family history and individuals must meet specific criteria, which are:
At least three relatives have colorectal or other HNPCC-related cancers, with one person being a first-degree relative (parent, child, or sibling) of the other two;
At least two generations have been affected;
At least one person received a diagnosis of colorectal cancer before age 50;
Familial adenomatous polyposis has been excluded."
maybe this is out of date? Is it usual  to have such briteria for a disease that can be described genetically?
Probably this additioanl information isn't necessary, but I think it would be useful to add some comment as to why these cannot be classified as HNPCC's because it isn't immediately obvius besed on my knowledge of the causative genes?
 in the definition
All of the colorectal cancer, hereditary nonpolyposis, should be  part of the colorectal cancer superclass, because these are the major hereditary cancer causing genes.
Does that sound reasonable?
@cmungall
@pnrobinson

These are the Amsterdam criteria and should be cited as such.
https://www.sciencedirect.com/topics/medicine-and-dentistry/amsterdam-criteria
The diagnosis of a disease should not be made based on a genetic variant alone because (i) it is hard to know if some variants are pathogenic and (ii) non-penetrance. Disease criteria should always have some clinical definitions!

From what I can tell, it seems like some sources use the terms Lynch syndrome and 'hereditary nonpolyposis colon cancer' synonymously, but NCIt has these as child and parent (as Val noted above).
Action items:
exclude superclass Lynch syndrome from the terms below. These classes should be children of MONDO_0018630 'hereditary nonpolyposis colon cancer'

[ ] 'colorectal cancer, hereditary nonpolyposis, type 2'

[ ] 'colorectal cancer, hereditary nonpolyposis, type 4'

[ ]  'colorectal cancer, hereditary nonpolyposis, type 5'

[ ] 'colorectal cancer, hereditary nonpolyposis, type 6'

[ ] 'colorectal cancer, hereditary nonpolyposis, type 7'

[ ] 'colorectal cancer, hereditary nonpolyposis, type 8'

[ ] 'Lynch syndrome 1'

[ ] change the xrefs to the OMIM classes on 'Lynch syndrome' from MONDO:superClassOf to MONDO:directSiblingOf

[ ] add comment to Lynch syndrome:
Sometimes Lynch syndrome is also referred to as hereditary non-polyposis colorectal cancer (HNPCC), but the two conditions are subtly different. Lynch syndrome is classified by a mutation in mismatch repair genes. It is diagnosed by specific criteria known as the Amsterdam criteria.

[ ] Add seeAlso on Lynch syndrome and 'hereditary nonpolyposis colon cancer': https://my.clevelandclinic.org/health/diseases/17195-lynch-syndrome--hnpcc

[ ] Add seeAlso to this link on Lynch syndrome: https://www.sciencedirect.com/topics/medicine-and-dentistry/amsterdam-criteria

[ ] Add seeAlso to this ticket on Lynch syndrome and 'hereditary nonpolyposis colon cancer'


[WIP] :construction: circles in skills graph

fix #187

According to the message of Julien Landuré (https://zenika.slack.com/archives/C021ASG3K8V/p1660657892218849) on the Skillz Slack. I purpose something like that :

And then when the mouse is hover a circle :

This app inspired me. And I think, it's a good idea to remove all the skills name that have the same note. Cause, if the user have too much skills, the circles can overlap, and we wouldn't see anythin
Please tell me what do you think @StephanieHP @benjaminbrgs

I agree with that. It's really better like this 👍🏼
I was thinking about having a modal when you click on the circle instead of zooming ? In this way, we could improve the readability of the skills (in a list for example).
Your solution is ok for me, mine is just a thinking and could be lead to too many modal.

Thanks @benjaminbrgs for the fix of the bigger circles :100:
I'll just fix something that I've missed :

It's showing the lenght on the home radar, but it should display the labels