From 297241e72286fc4154fe82063b8e320c7bdd1bba Mon Sep 17 00:00:00 2001
From: srinivamd <52507740+srinivamd@users.noreply.github.com>
Date: Thu, 21 May 2026 00:17:48 -0700
Subject: [PATCH] Fix 'Memory access fault' for fused dense on MI350 (#325)

Use getCurrentCUDABlasLtHandle() instead of getCurrentCUDABlasHandle()
in gemm_lt() to get the correct hipBLASLt handle, fixing a
write-to-read-only-page fault on gfx950.

Cherry-pick of af25af4e26f0546d897c3e2ff31502ef50530ecc
---
 csrc/fused_dense_cuda.cu | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/csrc/fused_dense_cuda.cu b/csrc/fused_dense_cuda.cu
index 15c076f68..50299c90d 100644
--- a/csrc/fused_dense_cuda.cu
+++ b/csrc/fused_dense_cuda.cu
@@ -142,7 +142,7 @@ int gemm_lt(
 {
 
   hipStream_t stream;
-  hipblasHandle_t handle = at::cuda::getCurrentCUDABlasHandle();
+  hipblasHandle_t handle = at::cuda::getCurrentCUDABlasLtHandle();
   hipblasGetStream(handle, &stream);
 
 #if DEBUG