1if (NOT EXISTS $ENV{ROCM_PATH})
  2    if (NOT EXISTS /opt/rocm)
  3        set(ROCM_PATH /usr)
  4    else()
  5        set(ROCM_PATH /opt/rocm)
  6    endif()
  7else()
  8    set(ROCM_PATH $ENV{ROCM_PATH})
  9endif()
 10
 11list(APPEND CMAKE_PREFIX_PATH  ${ROCM_PATH})
 12list(APPEND CMAKE_PREFIX_PATH "${ROCM_PATH}/lib64/cmake")
 13
 14# CMake on Windows doesn't support the HIP language yet
 15if (WIN32)
 16    set(CXX_IS_HIPCC TRUE)
 17else()
 18    string(REGEX MATCH "hipcc(\.bat)?$" CXX_IS_HIPCC "${CMAKE_CXX_COMPILER}")
 19endif()
 20
 21if (CXX_IS_HIPCC)
 22    if (LINUX)
 23        if (NOT ${CMAKE_CXX_COMPILER_ID} MATCHES "Clang")
 24            message(WARNING "Only LLVM is supported for HIP, hint: CXX=/opt/rocm/llvm/bin/clang++")
 25        endif()
 26
 27        message(WARNING "Setting hipcc as the C++ compiler is legacy behavior."
 28                " Prefer setting the HIP compiler directly. See README for details.")
 29    endif()
 30else()
 31    # Forward (AMD)GPU_TARGETS to CMAKE_HIP_ARCHITECTURES.
 32    if(AMDGPU_TARGETS AND NOT GPU_TARGETS)
 33        set(GPU_TARGETS ${AMDGPU_TARGETS})
 34    endif()
 35    if(GPU_TARGETS AND NOT CMAKE_HIP_ARCHITECTURES)
 36        set(CMAKE_HIP_ARCHITECTURES ${GPU_TARGETS})
 37    endif()
 38    cmake_minimum_required(VERSION 3.21)
 39    enable_language(HIP)
 40endif()
 41
 42find_package(hip     REQUIRED)
 43find_package(hipblas REQUIRED)
 44find_package(rocblas REQUIRED)
 45
 46if (${hip_VERSION} VERSION_LESS 6.1)
 47    message(FATAL_ERROR "At least ROCM/HIP V6.1 is required")
 48endif()
 49
 50message(STATUS "HIP and hipBLAS found")
 51
 52# Workaround old compilers
 53set(CMAKE_HIP_FLAGS "${CMAKE_HIP_FLAGS} --gpu-max-threads-per-block=1024")
 54
 55file(GLOB   GGML_HEADERS_ROCM "../ggml-cuda/*.cuh")
 56list(APPEND GGML_HEADERS_ROCM "../../include/ggml-cuda.h")
 57
 58file(GLOB   GGML_SOURCES_ROCM "../ggml-cuda/*.cu")
 59file(GLOB   SRCS "../ggml-cuda/template-instances/fattn-tile*.cu")
 60list(APPEND GGML_SOURCES_ROCM ${SRCS})
 61file(GLOB   SRCS "../ggml-cuda/template-instances/fattn-mma*.cu")
 62list(APPEND GGML_SOURCES_ROCM ${SRCS})
 63file(GLOB   SRCS "../ggml-cuda/template-instances/mmq*.cu")
 64list(APPEND GGML_SOURCES_ROCM ${SRCS})
 65file(GLOB   SRCS "../ggml-cuda/template-instances/mmf*.cu")
 66list(APPEND GGML_SOURCES_ROCM ${SRCS})
 67
 68if (GGML_CUDA_FA_ALL_QUANTS)
 69    file(GLOB   SRCS "../ggml-cuda/template-instances/fattn-vec*.cu")
 70    list(APPEND GGML_SOURCES_ROCM ${SRCS})
 71    add_compile_definitions(GGML_CUDA_FA_ALL_QUANTS)
 72else()
 73    file(GLOB   SRCS "../ggml-cuda/template-instances/fattn-vec*q4_0-q4_0.cu")
 74    list(APPEND GGML_SOURCES_ROCM ${SRCS})
 75    file(GLOB   SRCS "../ggml-cuda/template-instances/fattn-vec*q8_0-q8_0.cu")
 76    list(APPEND GGML_SOURCES_ROCM ${SRCS})
 77    file(GLOB   SRCS "../ggml-cuda/template-instances/fattn-vec*f16-f16.cu")
 78    list(APPEND GGML_SOURCES_ROCM ${SRCS})
 79endif()
 80
 81ggml_add_backend_library(ggml-hip
 82                         ${GGML_HEADERS_ROCM}
 83                         ${GGML_SOURCES_ROCM}
 84                        )
 85
 86# TODO: do not use CUDA definitions for HIP
 87if (NOT GGML_BACKEND_DL)
 88    target_compile_definitions(ggml PUBLIC GGML_USE_CUDA)
 89endif()
 90
 91add_compile_definitions(GGML_USE_HIP)
 92
 93if (GGML_CUDA_FORCE_MMQ)
 94    add_compile_definitions(GGML_CUDA_FORCE_MMQ)
 95endif()
 96
 97if (GGML_CUDA_FORCE_CUBLAS)
 98    add_compile_definitions(GGML_CUDA_FORCE_CUBLAS)
 99endif()
100
101if (GGML_CUDA_NO_PEER_COPY)
102    add_compile_definitions(GGML_CUDA_NO_PEER_COPY)
103endif()
104
105if (GGML_HIP_GRAPHS)
106    add_compile_definitions(GGML_HIP_GRAPHS)
107endif()
108
109if (GGML_HIP_NO_VMM)
110    add_compile_definitions(GGML_HIP_NO_VMM)
111endif()
112
113if (GGML_HIP_ROCWMMA_FATTN)
114    add_compile_definitions(GGML_HIP_ROCWMMA_FATTN)
115endif()
116
117if (NOT GGML_HIP_MMQ_MFMA)
118    add_compile_definitions(GGML_HIP_NO_MMQ_MFMA)
119endif()
120
121if (GGML_HIP_EXPORT_METRICS)
122    set(CMAKE_HIP_FLAGS "${CMAKE_HIP_FLAGS} -Rpass-analysis=kernel-resource-usage --save-temps")
123endif()
124
125if (NOT GGML_CUDA_FA)
126    add_compile_definitions(GGML_CUDA_NO_FA)
127endif()
128
129if (CXX_IS_HIPCC)
130    set_source_files_properties(${GGML_SOURCES_ROCM} PROPERTIES LANGUAGE CXX)
131    target_link_libraries(ggml-hip PRIVATE hip::device)
132else()
133    set_source_files_properties(${GGML_SOURCES_ROCM} PROPERTIES LANGUAGE HIP)
134endif()
135
136if (GGML_STATIC)
137    message(FATAL_ERROR "Static linking not supported for HIP/ROCm")
138endif()
139
140target_link_libraries(ggml-hip PRIVATE ggml-base hip::host roc::rocblas roc::hipblas)