mirror of
https://github.com/gentoo-mirror/guru.git
synced 2026-11-05 23:30:19 -05:00
Sv. Lockal
907f84730c
- curl is not a valid dependency in llama.cpp - openmp flags only wprk with clang, not hipcc - add patch for rocm-10.1 Signed-off-by: Sv. Lockal <lockalsash@gmail.com>
45 lines
2.6 KiB
Diff
45 lines
2.6 KiB
Diff
Fix build with LLVM 23
|
|
|
|
Since https://github.com/llvm/llvm-project/pull/176033 (LLVM 23), AMDGPU
|
|
builtins take ext_vector_type operands, which allow no implicit conversion
|
|
between element types. Older clang (e. g. 22 in ROCm 7.x) accepts both.
|
|
|
|
This is a short version of https://github.com/ggml-org/llama.cpp/pull/29536,
|
|
which also adds CDNA fix.
|
|
--- a/ggml/src/ggml-cuda/mma.cuh
|
|
+++ b/ggml/src/ggml-cuda/mma.cuh
|
|
@@ -1409,14 +1409,14 @@ namespace ggml_cuda_mma {
|
|
tile<16, 16, float, dl_d> & D, const tile<16, 8, nv_bfloat162, dl_ab> & A, const tile<16, 8, nv_bfloat162, dl_ab> & B) {
|
|
#if defined(AMD_WMMA_AVAILABLE)
|
|
#if defined(RDNA4)
|
|
- using bf16x8_t = __attribute__((ext_vector_type(8))) __bf16;
|
|
+ using bf16x8_t = __attribute__((ext_vector_type(8))) short;
|
|
using floatx8_t = __attribute__((ext_vector_type(8))) float;
|
|
floatx8_t& acc_frag = reinterpret_cast<floatx8_t&>(D.x[0]);
|
|
const bf16x8_t& a_frag = reinterpret_cast<const bf16x8_t&>(A.x[0]);
|
|
const bf16x8_t& b_frag = reinterpret_cast<const bf16x8_t&>(B.x[0]);
|
|
acc_frag = __builtin_amdgcn_wmma_f32_16x16x16_bf16_w32_gfx12(a_frag, b_frag, acc_frag);
|
|
#elif defined(RDNA3)
|
|
- using bf16x16_t = __attribute__((ext_vector_type(16))) __bf16;
|
|
+ using bf16x16_t = __attribute__((ext_vector_type(16))) short;
|
|
using floatx8_t = __attribute__((ext_vector_type(8))) float;
|
|
floatx8_t& acc_frag = reinterpret_cast<floatx8_t&>(D.x[0]);
|
|
const bf16x16_t& a_frag = reinterpret_cast<const bf16x16_t&>(A.x[0]);
|
|
@@ -1430,14 +1430,14 @@ namespace ggml_cuda_mma {
|
|
using floatx4_t = __attribute__((ext_vector_type(4))) float;
|
|
floatx4_t& acc_frag = reinterpret_cast<floatx4_t&>(D.x[0]);
|
|
#if defined(CDNA4) || defined(CDNA3) || defined(CDNA2)
|
|
- using bf16x4_t = __attribute__((ext_vector_type(4))) __bf16;
|
|
+ using bf16x4_t = __attribute__((ext_vector_type(4))) short;
|
|
const bf16x4_t& a_frag = reinterpret_cast<const bf16x4_t&>(A.x[0]);
|
|
const bf16x4_t& b_frag = reinterpret_cast<const bf16x4_t&>(B.x[0]);
|
|
acc_frag = __builtin_amdgcn_mfma_f32_16x16x16bf16_1k(a_frag, b_frag, acc_frag, 0, 0, 0);
|
|
#elif defined(CDNA1)
|
|
#pragma unroll
|
|
for (int i = 0; i < 2; ++i) {
|
|
- using bf16x2_t = __attribute__((ext_vector_type(2))) __bf16;
|
|
+ using bf16x2_t = __attribute__((ext_vector_type(2))) short;
|
|
const bf16x2_t& a_frag = reinterpret_cast<const bf16x2_t&>(A.x[i]);
|
|
const bf16x2_t& b_frag = reinterpret_cast<const bf16x2_t&>(B.x[i]);
|
|
acc_frag = __builtin_amdgcn_mfma_f32_16x16x8bf16(a_frag, b_frag, acc_frag, 0, 0, 0);
|