Skip to content

Commit aa79524

Browse files
IMbackKIMbackK
authored
HIP: remove the use of __HIP_PLATFORM_AMD__, explicitly support only AMD targets (ggml-org#14945)
1 parent b77d111 commit aa79524

File tree

10 files changed

+67
-65
lines changed

10 files changed

+67
-65
lines changed

‎ggml/src/ggml-cuda/common.cuh‎

Lines changed: 25 additions & 25 deletions
Original file line numberDiff line numberDiff line change
@@ -176,7 +176,7 @@ static const char * cu_get_error_str(CUresult err) {
176176
#define CU_CHECK(err) CUDA_CHECK_GEN(err, CUDA_SUCCESS, cu_get_error_str)
177177
#endif
178178

179-
#if !(defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)) && !defined(GGML_USE_MUSA)
179+
#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA)
180180
# define CUDA_SET_SHARED_MEMORY_LIMIT(kernel, nbytes) \
181181
do { \
182182
static bool shared_memory_limit_raised[GGML_CUDA_MAX_DEVICES] = { false }; \
@@ -191,7 +191,7 @@ static const char * cu_get_error_str(CUresult err) {
191191
do { \
192192
GGML_UNUSED(nbytes); \
193193
} while (0)
194-
#endif // !(defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)) && !defined(GGML_USE_MUSA)
194+
#endif // !(defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA)
195195

196196
#if CUDART_VERSION >= 11010 || defined(GGML_USE_MUSA)
197197
#define GGML_CUDA_ASSUME(x) __builtin_assume(x)
@@ -211,9 +211,9 @@ typedef float2 dfloat2;
211211
#define GGML_USE_VMM
212212
#endif // (!defined(GGML_USE_HIP) && !defined(GGML_CUDA_NO_VMM)) || (defined(GGML_USE_HIP) && !defined(GGML_HIP_NO_VMM))
213213

214-
#if (defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)) || __CUDA_ARCH__ >= GGML_CUDA_CC_PASCAL
214+
#if defined(GGML_USE_HIP) || __CUDA_ARCH__ >= GGML_CUDA_CC_PASCAL
215215
#define FP16_AVAILABLE
216-
#endif // (defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)) || __CUDA_ARCH__ >= GGML_CUDA_CC_PASCAL
216+
#endif // defined(GGML_USE_HIP) || __CUDA_ARCH__ >= GGML_CUDA_CC_PASCAL
217217

218218
#if defined(FP16_AVAILABLE) && __CUDA_ARCH__ != 610
219219
#define FAST_FP16_AVAILABLE
@@ -227,17 +227,17 @@ typedef float2 dfloat2;
227227
#define FP16_MMA_AVAILABLE
228228
#endif // defined(GGML_HIP_ROCWMMA_FATTN) && (defined(CDNA) || defined(RDNA3) || (defined(GGML_HIP_ROCWMMA_FATTN_GFX12) && defined(RDNA4)))
229229

230-
#if defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__) && defined(CDNA3) && !defined(GGML_HIP_NO_MMQ_MFMA)
230+
#if defined(GGML_USE_HIP) && defined(CDNA3) && !defined(GGML_HIP_NO_MMQ_MFMA)
231231
#define AMD_MFMA_AVAILABLE
232-
#endif // defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__) && defined(CDNA3)
232+
#endif // defined(GGML_USE_HIP) && defined(CDNA3) && !defined(GGML_HIP_NO_MMQ_MFMA)
233233

234-
#if !(defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)) && __CUDA_ARCH__ >= GGML_CUDA_CC_TURING
234+
#if !defined(GGML_USE_HIP) && __CUDA_ARCH__ >= GGML_CUDA_CC_TURING
235235
#define NEW_MMA_AVAILABLE
236-
#endif // !(defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)) && __CUDA_ARCH__ >= GGML_CUDA_CC_TURING
236+
#endif // !defined(GGML_USE_HIP) && __CUDA_ARCH__ >= GGML_CUDA_CC_TURING
237237

238-
#if !(defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)) && __CUDA_ARCH__ >= GGML_CUDA_CC_AMPERE
238+
#if !defined(GGML_USE_HIP) && __CUDA_ARCH__ >= GGML_CUDA_CC_AMPERE
239239
#define CP_ASYNC_AVAILABLE
240-
#endif // !(defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)) && __CUDA_ARCH__ >= GGML_CUDA_CC_AMPERE
240+
#endif // !defined(GGML_USE_HIP) && __CUDA_ARCH__ >= GGML_CUDA_CC_AMPERE
241241

242242
#if !defined(GGML_CUDA_NO_FA) && !(defined(GGML_USE_MUSA) && __MUSA_ARCH__ < 220)
243243
#define FLASH_ATTN_AVAILABLE
@@ -259,7 +259,7 @@ static bool fast_fp16_hardware_available(const int cc) {
259259

260260
// Any FP16 tensor core instructions are available for ggml code.
261261
static bool fp16_mma_available(const int cc) {
262-
#if defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__) && !defined(GGML_HIP_ROCWMMA_FATTN)
262+
#if defined(GGML_USE_HIP) && !defined(GGML_HIP_ROCWMMA_FATTN)
263263
return false;
264264
#else
265265
if ((GGML_CUDA_CC_IS_NVIDIA(cc) && ggml_cuda_highest_compiled_arch(cc) >= GGML_CUDA_CC_VOLTA) ||
@@ -275,7 +275,7 @@ static bool fp16_mma_available(const int cc) {
275275
} else {
276276
return false;
277277
}
278-
#endif // defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__) && !defined(GGML_HIP_ROCWMMA_FATTN)
278+
#endif // defined(GGML_USE_HIP) && !defined(GGML_HIP_ROCWMMA_FATTN)
279279
}
280280

281281
// To be used for feature selection of external libraries, e.g. cuBLAS.
@@ -312,25 +312,25 @@ static bool cp_async_available(const int cc) {
312312
}
313313

314314
static constexpr __device__ int ggml_cuda_get_physical_warp_size() {
315-
#if defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__) && (defined(__GFX9__) || defined(__GFX8__))
315+
#if defined(GGML_USE_HIP) && (defined(__GFX9__) || defined(__GFX8__))
316316
return 64;
317317
#else
318318
return 32;
319-
#endif // defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__) && (defined(__GFX9__) || defined(__GFX8__))
319+
#endif // defined(GGML_USE_HIP) && (defined(__GFX9__) || defined(__GFX8__))
320320
}
321321

322322
[[noreturn]]
323323
static __device__ void no_device_code(
324324
const char * file_name, const int line, const char * function_name, const int arch, const char * arch_list) {
325325

326-
#if defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)
326+
#if defined(GGML_USE_HIP)
327327
printf("%s:%d: ERROR: HIP kernel %s has no device code compatible with HIP arch %d.\n",
328328
file_name, line, function_name, arch);
329329
GGML_UNUSED(arch_list);
330330
#else
331331
printf("%s:%d: ERROR: CUDA kernel %s has no device code compatible with CUDA arch %d. ggml-cuda.cu was compiled for: %s\n",
332332
file_name, line, function_name, arch, arch_list);
333-
#endif // defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)
333+
#endif // defined(GGML_USE_HIP)
334334
__trap();
335335

336336
GGML_UNUSED(no_device_code); // suppress unused function warning
@@ -367,15 +367,15 @@ struct ggml_cuda_unroll<1> {
367367

368368
template<int width = WARP_SIZE>
369369
static __device__ __forceinline__ int warp_reduce_sum(int x) {
370-
#if !(defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)) && __CUDA_ARCH__ >= GGML_CUDA_CC_AMPERE
370+
#if !defined(GGML_USE_HIP) && __CUDA_ARCH__ >= GGML_CUDA_CC_AMPERE
371371
return __reduce_add_sync(0xffffffff, x);
372372
#else
373373
#pragma unroll
374374
for (int offset = width/2; offset > 0; offset >>= 1) {
375375
x += __shfl_xor_sync(0xffffffff, x, offset, width);
376376
}
377377
return x;
378-
#endif // !(defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)) && __CUDA_ARCH__ >= GGML_CUDA_CC_AMPERE
378+
#endif // !defined(GGML_USE_HIP) && __CUDA_ARCH__ >= GGML_CUDA_CC_AMPERE
379379
}
380380

381381
template<int width = WARP_SIZE>
@@ -444,11 +444,11 @@ static __device__ __forceinline__ float warp_reduce_max(float x) {
444444
static __device__ __forceinline__ half ggml_cuda_hmax(const half a, const half b) {
445445
#ifdef FP16_AVAILABLE
446446

447-
#if !(defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)) && CUDART_VERSION < CUDART_HMAX
447+
#if !defined(GGML_USE_HIP) && CUDART_VERSION < CUDART_HMAX
448448
return __float2half(fmaxf(__half2float(a), __half2float(b)));
449449
#else
450450
return __hmax(a, b);
451-
#endif // !(defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)) && CUDART_VERSION < CUDART_HMAX
451+
#endif // !defined(GGML_USE_HIP) && CUDART_VERSION < CUDART_HMAX
452452

453453
#else
454454
NO_DEVICE_CODE;
@@ -476,7 +476,7 @@ static __device__ __forceinline__ half2 ggml_cuda_hmax2(const half2 a, const hal
476476

477477
template<int width = WARP_SIZE>
478478
static __device__ __forceinline__ half2 warp_reduce_max(half2 x) {
479-
#if !(defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)) && __CUDA_ARCH__ >= GGML_CUDA_CC_PASCAL || (defined(GGML_USE_HIP) && HIP_VERSION >= 50700000)
479+
#if !defined(GGML_USE_HIP) && __CUDA_ARCH__ >= GGML_CUDA_CC_PASCAL || (defined(GGML_USE_HIP) && HIP_VERSION >= 50700000)
480480
#pragma unroll
481481
for (int offset = width/2; offset > 0; offset >>= 1) {
482482
x = ggml_cuda_hmax2(x, __shfl_xor_sync(0xffffffff, x, offset, width));
@@ -485,7 +485,7 @@ static __device__ __forceinline__ half2 warp_reduce_max(half2 x) {
485485
#else
486486
GGML_UNUSED(x);
487487
NO_DEVICE_CODE;
488-
#endif // !(defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)) && __CUDA_ARCH__ >= GGML_CUDA_CC_PASCAL || (defined(GGML_USE_HIP) && HIP_VERSION >= 50700000)
488+
#endif // !defined(GGML_USE_HIP) && __CUDA_ARCH__ >= GGML_CUDA_CC_PASCAL || (defined(GGML_USE_HIP) && HIP_VERSION >= 50700000)
489489
}
490490

491491
#if CUDART_VERSION < CUDART_HMASK
@@ -497,7 +497,7 @@ static __device__ __forceinline__ uint32_t __hgt2_mask(const half2 a, const half
497497
#endif // CUDART_VERSION < CUDART_HMASK
498498

499499
static __device__ __forceinline__ int ggml_cuda_dp4a(const int a, const int b, int c) {
500-
#if defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)
500+
#if defined(GGML_USE_HIP)
501501
#if defined(CDNA) || defined(RDNA2) || defined(__gfx906__)
502502
c = __builtin_amdgcn_sdot4(a, b, c, false);
503503
#elif defined(RDNA3) || defined(RDNA4)
@@ -523,7 +523,7 @@ static __device__ __forceinline__ int ggml_cuda_dp4a(const int a, const int b, i
523523
#endif
524524
return c;
525525

526-
#else // defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)
526+
#else // defined(GGML_USE_HIP)
527527

528528
#if __CUDA_ARCH__ >= GGML_CUDA_CC_DP4A || defined(GGML_USE_MUSA)
529529
return __dp4a(a, b, c);
@@ -533,7 +533,7 @@ static __device__ __forceinline__ int ggml_cuda_dp4a(const int a, const int b, i
533533
return c + a8[0]*b8[0] + a8[1]*b8[1] + a8[2]*b8[2] + a8[3]*b8[3];
534534
#endif // __CUDA_ARCH__ >= GGML_CUDA_CC_DP4A || defined(GGML_USE_MUSA)
535535

536-
#endif // defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)
536+
#endif // defined(GGML_USE_HIP)
537537
}
538538

539539
typedef void (*dequantize_kernel_t)(const void * vx, const int64_t ib, const int iqs, dfloat2 & v);

‎ggml/src/ggml-cuda/fattn-common.cuh‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -592,9 +592,9 @@ static __global__ void flash_attn_stream_k_fixup(
592592
}
593593

594594
template<int D> // D == head size
595-
#if !(defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__))
595+
#if !defined(GGML_USE_HIP)
596596
__launch_bounds__(D, 1)
597-
#endif // !(defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__))
597+
#endif // !(defined(GGML_USE_HIP)
598598
static __global__ void flash_attn_combine_results(
599599
const float * __restrict__ VKQ_parts,
600600
const float2 * __restrict__ VKQ_meta,

‎ggml/src/ggml-cuda/fattn-mma-f16.cuh‎

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -1391,24 +1391,24 @@ void ggml_cuda_flash_attn_ext_mma_f16_case(ggml_backend_cuda_context & ctx, ggml
13911391
constexpr bool use_logit_softcap = false;
13921392
fattn_kernel = flash_attn_ext_f16<DKQ, DV, ncols1, ncols2, nwarps, ntiles, use_logit_softcap, mla>;
13931393

1394-
#if !(defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)) && !defined(GGML_USE_MUSA)
1394+
#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA)
13951395
static bool shared_memory_limit_raised[GGML_CUDA_MAX_DEVICES] = {false};
13961396
if (!shared_memory_limit_raised[id]) {
13971397
CUDA_CHECK(cudaFuncSetAttribute(fattn_kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, nbytes_shared_total));
13981398
shared_memory_limit_raised[id] = true;
13991399
}
1400-
#endif // !(defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)) && !defined(GGML_USE_MUSA)
1400+
#endif // !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA)
14011401
} else {
14021402
constexpr bool use_logit_softcap = true;
14031403
fattn_kernel = flash_attn_ext_f16<DKQ, DV, ncols1, ncols2, nwarps, ntiles, use_logit_softcap, mla>;
14041404

1405-
#if !(defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)) && !defined(GGML_USE_MUSA)
1405+
#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA)
14061406
static bool shared_memory_limit_raised[GGML_CUDA_MAX_DEVICES] = {false};
14071407
if (!shared_memory_limit_raised[id]) {
14081408
CUDA_CHECK(cudaFuncSetAttribute(fattn_kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, nbytes_shared_total));
14091409
shared_memory_limit_raised[id] = true;
14101410
}
1411-
#endif // !(defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)) && !defined(GGML_USE_MUSA)
1411+
#endif // !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA)
14121412
}
14131413

14141414
launch_fattn<DV, ncols1, ncols2>

‎ggml/src/ggml-cuda/fattn-tile-f16.cu‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -5,9 +5,9 @@
55
#define FATTN_KQ_STRIDE_TILE_F16 64
66

77
template<int D, int ncols, int nwarps, bool use_logit_softcap> // D == head size
8-
#if !(defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__))
8+
#if !defined(GGML_USE_HIP)
99
__launch_bounds__(nwarps*WARP_SIZE, 2)
10-
#endif // !(defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__))
10+
#endif // !defined(GGML_USE_HIP)
1111
static __global__ void flash_attn_tile_ext_f16(
1212
const char * __restrict__ Q,
1313
const char * __restrict__ K,

‎ggml/src/ggml-cuda/fattn-tile-f32.cu‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -5,9 +5,9 @@
55
#define FATTN_KQ_STRIDE_TILE_F32 32
66

77
template<int D, int ncols, int nwarps, bool use_logit_softcap> // D == head size
8-
#if !(defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__))
8+
#if !defined(GGML_USE_HIP)
99
__launch_bounds__(nwarps*WARP_SIZE, 2)
10-
#endif // !(defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__))
10+
#endif // !defined(GGML_USE_HIP)
1111
static __global__ void flash_attn_tile_ext_f32(
1212
const char * __restrict__ Q,
1313
const char * __restrict__ K,

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

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -7,7 +7,7 @@
77
#include "fattn-wmma-f16.cuh"
88

99
#ifdef FP16_MMA_AVAILABLE
10-
#if !(defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__))
10+
#if !defined(GGML_USE_HIP)
1111
#include <mma.h>
1212
#ifdef GGML_USE_MUSA
1313
namespace wmma = mtmusa::wmma;
@@ -18,7 +18,7 @@ namespace wmma = nvcuda::wmma;
1818
#undef HIP_ENABLE_WARP_SYNC_BUILTINS // conflicts with rocWMMA headers
1919
#include <rocwmma/rocwmma.hpp>
2020
namespace wmma = rocwmma;
21-
#endif // !(defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__))
21+
#endif // !defined(GGML_USE_HIP)
2222
#endif // FP16_MMA_AVAILABLE
2323

2424
// D == head size, VKQ_stride == num VKQ rows calculated in parallel:
@@ -546,7 +546,7 @@ void ggml_cuda_flash_attn_ext_wmma_f16(ggml_backend_cuda_context & ctx, ggml_ten
546546
return;
547547
}
548548

549-
#if !(defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__))
549+
#if !defined(GGML_USE_HIP)
550550
if (Q->ne[1] <= 8 && Q->ne[0] % warp_size == 0) {
551551
constexpr int cols_per_block = 8;
552552
switch (Q->ne[0]) {
@@ -568,7 +568,7 @@ void ggml_cuda_flash_attn_ext_wmma_f16(ggml_backend_cuda_context & ctx, ggml_ten
568568
}
569569
return;
570570
}
571-
#endif // !(defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__))
571+
#endif // !defined(GGML_USE_HIP)
572572

573573
if (Q->ne[1] <= 32) {
574574
constexpr int cols_per_block = 16;

‎ggml/src/ggml-cuda/ggml-cuda.cu‎

Lines changed: 5 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -128,7 +128,7 @@ static cudaError_t ggml_cuda_device_malloc(void ** ptr, size_t size, int device)
128128
return err;
129129
}
130130

131-
#if defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)
131+
#if defined(GGML_USE_HIP)
132132
static int ggml_cuda_parse_id(char devName[]) {
133133
// A list of possible Target IDs can be found under the rocclr/clr repo in device.cpp
134134
// these values are not stable so this is susceptible to breakage
@@ -175,10 +175,10 @@ static int ggml_cuda_parse_id(char devName[]) {
175175
archNum += archMinor;
176176
return archNum;
177177
}
178-
#endif // defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)
178+
#endif // defined(GGML_USE_HIP)
179179

180180
static ggml_cuda_device_info ggml_cuda_init() {
181-
#ifdef __HIP_PLATFORM_AMD__
181+
#if defined(GGML_USE_HIP)
182182
// Workaround for a rocBLAS bug when using multiple graphics cards:
183183
// https://github.com/ROCmSoftwarePlatform/rocBLAS/issues/1346
184184
{
@@ -251,7 +251,7 @@ static ggml_cuda_device_info ggml_cuda_init() {
251251
info.devices[id].nsm = prop.multiProcessorCount;
252252
info.devices[id].smpb = prop.sharedMemPerBlock;
253253
info.devices[id].warp_size = prop.warpSize;
254-
#if defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)
254+
#if defined(GGML_USE_HIP)
255255
info.devices[id].smpbo = prop.sharedMemPerBlock;
256256

257257
info.devices[id].cc = ggml_cuda_parse_id(prop.gcnArchName);
@@ -281,7 +281,7 @@ static ggml_cuda_device_info ggml_cuda_init() {
281281
info.devices[id].cc = 100*prop.major + 10*prop.minor;
282282
GGML_LOG_INFO(" Device %d: %s, compute capability %d.%d, VMM: %s\n",
283283
id, prop.name, prop.major, prop.minor, device_vmm ? "yes" : "no");
284-
#endif // defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)
284+
#endif // defined(GGML_USE_HIP)
285285
}
286286

287287
for (int id = 0; id < info.device_count; ++id) {

‎ggml/src/ggml-cuda/mma.cuh‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -68,7 +68,7 @@ namespace ggml_cuda_mma {
6868
static constexpr int I = I_;
6969
static constexpr int J = J_;
7070

71-
#if defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)
71+
#if defined(GGML_USE_HIP)
7272
static constexpr int ne = I * J / 64;
7373
T x[ne] = {0};
7474

@@ -132,7 +132,7 @@ namespace ggml_cuda_mma {
132132
static_assert(I == -1 && J == -1, "template specialization not implemented");
133133
}
134134
}
135-
#endif // defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)
135+
#endif // defined(GGML_USE_HIP)
136136
};
137137

138138
template <int I_, int J_>

0 commit comments

Comments
 (0)