Skip to content

Commit 7429c42

Browse files
committed
Support gfx*-generic targets
1 parent acdc3f0 commit 7429c42

File tree

3 files changed

+21
-4
lines changed

3 files changed

+21
-4
lines changed

ggml/src/ggml-cuda/fattn-wmma-f16.cu

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -15,6 +15,12 @@ namespace wmma = mtmusa::wmma;
1515
namespace wmma = nvcuda::wmma;
1616
#endif // GGML_USE_MUSA
1717
#elif defined(GGML_HIP_ROCWMMA_FATTN) && defined(FP16_MMA_AVAILABLE)
18+
// Workaround for gfx*-generic
19+
#if defined(__gfx11_generic__)
20+
#define __gfx1100__ __gfx11_generic__
21+
#elif defined(__gfx12_generic__)
22+
#define __gfx1201__ __gfx12_generic__
23+
#endif
1824
#undef HIP_ENABLE_WARP_SYNC_BUILTINS // conflicts with rocWMMA headers
1925
#include <rocwmma/rocwmma.hpp>
2026
namespace wmma = rocwmma;

ggml/src/ggml-cuda/sum.cu

Lines changed: 11 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -5,6 +5,17 @@
55
#ifdef USE_CUB
66

77
#if defined(GGML_USE_HIP)
8+
// Workaround for gfx*-generic
9+
#if defined(__gfx10_1_generic__)
10+
#define __gfx1010__ __gfx10_1_generic__
11+
#elif defined(__gfx10_3_generic__)
12+
#define __gfx1030__ __gfx10_3_generic__
13+
#elif defined(__gfx11_generic__)
14+
#define __gfx1100__ __gfx11_generic__
15+
#elif defined(__gfx12_generic__)
16+
#define __gfx1201__ __gfx12_generic__
17+
#endif
18+
819
#include <hipcub/hipcub.hpp>
920
using namespace hipcub;
1021
#else

ggml/src/ggml-cuda/vendors/hip.h

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -158,17 +158,17 @@
158158
#define RDNA4
159159
#endif
160160

161-
#if defined(__gfx1100__) || defined(__gfx1101__) || defined(__gfx1102__) || defined(__gfx1103__) || \
162-
defined(__gfx1150__) || defined(__gfx1151__)
161+
#if defined(__GFX11__)
163162
#define RDNA3
164163
#endif
165164

166165
#if defined(__gfx1030__) || defined(__gfx1031__) || defined(__gfx1032__) || defined(__gfx1033__) || \
167-
defined(__gfx1034__) || defined(__gfx1035__) || defined(__gfx1036__) || defined(__gfx1037__)
166+
defined(__gfx1034__) || defined(__gfx1035__) || defined(__gfx1036__) || defined(__gfx1037__) || \
167+
defined(__gfx10_3_generic__)
168168
#define RDNA2
169169
#endif
170170

171-
#if defined(__gfx1010__) || defined(__gfx1012__)
171+
#if defined(__gfx1010__) || defined(__gfx1012__) || defined(__gfx10_1_generic__)
172172
#define RDNA1
173173
#endif
174174

0 commit comments

Comments
 (0)