Niko1221/Strata
1
0 Fork 0
mirror of https://github.com/Niko1221/Strata.git synced 2026-11-07 18:35:04 -10:00
Code Issues Packages Projects Releases 1 Wiki Activity

hip: the Q2_0 signed-zero fix also on gfx1200/gfx1201 (ROCm 7.10), where hip_q2_zero failed (#1474)

Browse Source
This commit is contained in:
Niko1221 committed 2026-10-09 03:15:13 +02:00
1 parent 3d6793a408
commit 0ff45cffb2
1 file changed
+4 -3
Download Patch File Download Diff File Expand all files Collapse all files
src/kernels/cuda/iq_kernels.cu
+4 -3
View File
@@ -1708,9 +1708,10 @@ __device__ void dq_q2_0(const void* vx, int64_t ibs, dst_t* yy, int tid) {
for (int j = 0; j < 8; ++j) {
const int i = part * 8 + j;
const int code = (x[b].qs[i / 4] >> ((i % 4) * 2)) & 3;
#if defined(__HIPCC__) && ((defined(__gfx1012__) && HIP_VERSION_MAJOR < 7) || defined(__gfx1151__))
// HIP 5.7 on RDNA1 and the HIP compiler on gfx1151 (PR #895, ROCm 7.13) fold the half path's negative scale
// times +0 to +0. Preserve the scale's sign, as the FP32/CPU paths do.
#if defined(__HIPCC__) && ((defined(__gfx1012__) && HIP_VERSION_MAJOR < 7) || defined(__gfx1151__) || \
defined(__gfx1200__) || defined(__gfx1201__))
// HIP 5.7 on RDNA1, the HIP compiler on gfx1151 (PR #895, ROCm 7.13) and on RDNA4 gfx1200/gfx1201 (ROCm
// 7.10, #1474) fold the half path's negative scale times +0 to +0. Preserve the scale's sign, as the FP32/CPU paths do.
if constexpr (std::is_same_v<dst_t, __half>) {
if (code == 1) {
yy[b * 64 + i] = __ushort_as_half(__half_as_ushort(x[b].d) & 0x8000u);