fix hip build

slaren · slaren · commit 1ea3769f07bf · 2025-08-05T18:39:14.000+02:00
ggml-ci
diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh
@@ -754,9 +754,9 @@ template <int mmq_y, bool need_check> static __device__ __forceinline__ void loa
         const block_mxfp4 * bxi = (const block_mxfp4 *) x + kbx0 + i*stride + kbxd;
 
 #if defined(AMD_MFMA_AVAILABLE) || defined(NEW_MMA_AVAILABLE)
-        x_df[i*MMQ_MMA_TILE_X_K_Q8_1                 + kbxd] = __bfloat162float(ggml_cuda_e8m0_to_fp32(bxi->e))*0.5f;
+        x_df[i*MMQ_MMA_TILE_X_K_Q8_1                 + kbxd] = ggml_cuda_e8m0_to_fp32(bxi->e)*0.5f;
 #else
-        x_df[i*(MMQ_TILE_NE_K/QI_MXFP4) + i/QI_MXFP4 + kbxd] = __bfloat162float(ggml_cuda_e8m0_to_fp32(bxi->e))*0.5f;
+        x_df[i*(MMQ_TILE_NE_K/QI_MXFP4) + i/QI_MXFP4 + kbxd] = ggml_cuda_e8m0_to_fp32(bxi->e)*0.5f;
 #endif // defined(AMD_MFMA_AVAILABLE) || defined(NEW_MMA_AVAILABLE)
     }
 }

Original file line number	Diff line number	Diff line change
`@@ -754,9 +754,9 @@ template <int mmq_y, bool need_check> static __device__ __forceinline__ void loa`
`754`	`754`	`const block_mxfp4 * bxi = (const block_mxfp4 ) x + kbx0 + istride + kbxd;`
`755`	`755`
`756`	`756`	`#if defined(AMD_MFMA_AVAILABLE) \|\| defined(NEW_MMA_AVAILABLE)`
`757`		`- x_df[iMMQ_MMA_TILE_X_K_Q8_1 + kbxd] = __bfloat162float(ggml_cuda_e8m0_to_fp32(bxi->e))0.5f;`
	`757`	`+ x_df[iMMQ_MMA_TILE_X_K_Q8_1 + kbxd] = ggml_cuda_e8m0_to_fp32(bxi->e)0.5f;`
`758`	`758`	`#else`
`759`		`- x_df[i(MMQ_TILE_NE_K/QI_MXFP4) + i/QI_MXFP4 + kbxd] = __bfloat162float(ggml_cuda_e8m0_to_fp32(bxi->e))0.5f;`
	`759`	`+ x_df[i(MMQ_TILE_NE_K/QI_MXFP4) + i/QI_MXFP4 + kbxd] = ggml_cuda_e8m0_to_fp32(bxi->e)0.5f;`
`760`	`760`	`#endif // defined(AMD_MFMA_AVAILABLE) \|\| defined(NEW_MMA_AVAILABLE)`
`761`	`761`	`}`
`762`	`762`	`}`