fix(kvarn): cast HIP kernel pointer for dynamic shmem attribute

Author: Anbeeld

ggml/src/ggml-cuda/kvarn.cu

@@ -638,8 +638,16 @@ void ggml_cuda_op_kvarn_store(ggml_backend_cuda_context & ctx, ggml_tensor * dst
     const size_t smpbo = ggml_cuda_info().devices[ctx.device].smpbo;
 
     if (!force_low_shmem && smpbo >= KVAR_N_SHARED_BYTES) {
-#if !defined(GGML_USE_MUSA)
-        CUDA_CHECK(cudaFuncSetAttribute(reinterpret_cast<const void*>(kvarn_store_kernel_hishmem), cudaFuncAttributeMaxDynamicSharedMemorySize, KVAR_N_SHARED_BYTES));
+#if defined(GGML_USE_HIP)
+        CUDA_CHECK(hipFuncSetAttribute(
+            reinterpret_cast<const void *>(&kvarn_store_kernel_hishmem),
+            hipFuncAttributeMaxDynamicSharedMemorySize,
+            KVAR_N_SHARED_BYTES));
+#elif !defined(GGML_USE_MUSA)
+        CUDA_CHECK(cudaFuncSetAttribute(
+            kvarn_store_kernel_hishmem,
+            cudaFuncAttributeMaxDynamicSharedMemorySize,
+            KVAR_N_SHARED_BYTES));
 #endif
         kvarn_store_kernel_hishmem<<<current->ne[1], KVAR_N_DIM, KVAR_N_SHARED_BYTES, ctx.stream()>>>(
             (const float *) current->data,