1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
|
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);
|
