Skip to content

Commit bb4206d

Browse files
committed
Ext: Fix register call for float
ROCm doesn't support the register keyword. This should fix compile. Signed-off-by: kingbri <8082010+kingbri1@users.noreply.github.com>
1 parent 0a77331 commit bb4206d

File tree

1 file changed

+4
-0
lines changed

1 file changed

+4
-0
lines changed

exllamav2/exllamav2_ext/cuda/q_mlp_softmax.cuh

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -279,7 +279,11 @@ __global__ void softmax128_topk_norm_kernel
279279
const int row = blockIdx.y * WARPSIZE + threadIdx.x;
280280
if (row >= rows) return;
281281

282+
#if defined(USE_ROCM)
283+
float f[128];
284+
#else
282285
register float f[128];
286+
#endif
283287

284288
int4* row_ptr = reinterpret_cast<int4*>(x + row * 128);
285289

0 commit comments

Comments
 (0)