Skip to content

Commit 68428ee

Browse files
committed
fix: add reinterpret_cast for hipFuncSetAttribute in kvarn ROCm build
HIP's hipFuncSetAttribute requires const void* but cannot implicitly convert a kernel function pointer, unlike native CUDA. Use reinterpret_cast<const void*> to match the pattern already used in fattn-mma-turbo.cuh and fattn-mma-f16.cuh.
1 parent 5ec9c50 commit 68428ee

1 file changed

Lines changed: 1 addition & 1 deletion

File tree

ggml/src/ggml-cuda/kvarn.cu

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -639,7 +639,7 @@ void ggml_cuda_op_kvarn_store(ggml_backend_cuda_context & ctx, ggml_tensor * dst
639639

640640
if (!force_low_shmem && smpbo >= KVAR_N_SHARED_BYTES) {
641641
#if !defined(GGML_USE_MUSA)
642-
CUDA_CHECK(cudaFuncSetAttribute(kvarn_store_kernel_hishmem, cudaFuncAttributeMaxDynamicSharedMemorySize, KVAR_N_SHARED_BYTES));
642+
CUDA_CHECK(cudaFuncSetAttribute(reinterpret_cast<const void*>(kvarn_store_kernel_hishmem), cudaFuncAttributeMaxDynamicSharedMemorySize, KVAR_N_SHARED_BYTES));
643643
#endif
644644
kvarn_store_kernel_hishmem<<<current->ne[1], KVAR_N_DIM, KVAR_N_SHARED_BYTES, ctx.stream()>>>(
645645
(const float *) current->data,

0 commit comments

Comments
 (0)