From 143c9075acb4b8630857ed2b468d657eb7c476c6 Mon Sep 17 00:00:00 2001 From: MISHANMAURYA <118961433+MISHANMAURYA@users.noreply.github.com> Date: Wed, 11 Jun 2025 22:36:41 +0530 Subject: [PATCH 01/12] Update ops_hip.cuh --- csrc/ops_hip.cuh | 8 +++++--- 1 file changed, 5 insertions(+), 3 deletions(-) diff --git a/csrc/ops_hip.cuh b/csrc/ops_hip.cuh index bcfc73e99..ff5c8e6ce 100644 --- a/csrc/ops_hip.cuh +++ b/csrc/ops_hip.cuh @@ -99,17 +99,18 @@ typedef enum Funcs_t class Context { public: - rocblas_handle m_handle; + hipblasLtHandle_t m_handle; Context() { - rocblas_handle handle; - rocblas_create_handle(&handle); + hipblasLtHandle_t handle; + hipblasLtCreate(&handle); m_handle = handle; } }; +/* class ContextLt { public: @@ -122,6 +123,7 @@ class ContextLt m_handle = handle; } }; +*/ class ContextHipsparse { From 1ee8f5692f6904e76710d87a3bae3f2c5c6427d6 Mon Sep 17 00:00:00 2001 From: MISHANMAURYA <118961433+MISHANMAURYA@users.noreply.github.com> Date: Fri, 13 Jun 2025 15:06:17 +0530 Subject: [PATCH 02/12] Update ops.hip --- csrc/ops.hip | 95 ++++++++++++++++++++++++++++++++++++++++++++++++++-- 1 file changed, 92 insertions(+), 3 deletions(-) diff --git a/csrc/ops.hip b/csrc/ops.hip index eef616d48..93b27ed8e 100644 --- a/csrc/ops.hip +++ b/csrc/ops.hip @@ -476,13 +476,51 @@ template int igemmlt( hipblasLtPointerMode_t pointerMode = HIPBLASLT_POINTER_MODE_ALPHA_DEVICE_VECTOR_BETA_HOST; has_error |= checkHipblasStatus(hipblasLtMatrixLayoutCreate(&aDesc, HIP_R_8I, m, k, lda)); + if(has_error != 0) + { + std::cout<<"failed to run hipblasLtMatrixLayoutCreate for Adesc:"< int igemmlt( fprintf(stderr, "Error: Matmul Algo Heuristic didn't return algorithms\n"); } else { int alpha = 1, beta = 0; + void* d_workspace = nullptr; + uint64_t workspace_size = 0; + for(int i = 0; i < returnedAlgoCount; i++) + workspace_size = max(workspace_size, heuristicResult[i].workspaceSize); + hipMalloc(&d_workspace, workspace_size); has_error |= checkHipblasStatus(hipblasLtMatmul( ltHandle, matmulDesc, &alpha, A, aDesc, B, bDesc, &beta, (int32_t*)C, cDesc, (int32_t*)C, cDesc, - &heuristicResult[0].algo, NULL, 0, stream + &heuristicResult[0].algo, d_workspace, workspace_size, stream )); + hipFree(d_workspace); + if(has_error != 0) + { + std::cout<<"failed to run hipblasLtMatmul"< int igemmlt( &pointerMode, sizeof(alphaVec) )); + if(has_error != 0) + { + std::cout<<"failed to run hipblasLtMatmulDescSetAttribute HIPBLASLT_MATMUL_DESC_POINTER_MODE for int8"< Date: Fri, 13 Jun 2025 15:16:28 +0530 Subject: [PATCH 03/12] Update ops_hip.cuh --- csrc/ops_hip.cuh | 3 +++ 1 file changed, 3 insertions(+) diff --git a/csrc/ops_hip.cuh b/csrc/ops_hip.cuh index ff5c8e6ce..80abe3595 100644 --- a/csrc/ops_hip.cuh +++ b/csrc/ops_hip.cuh @@ -100,9 +100,12 @@ class Context { public: hipblasLtHandle_t m_handle; + //rocblas_handle m_handle; Context() { + //rocblas_handle handle; + //rocblas_create_handle(&handle); hipblasLtHandle_t handle; hipblasLtCreate(&handle); m_handle = handle; From e9682e213cac38f9ef59c85405b95f0cfdc87ce8 Mon Sep 17 00:00:00 2001 From: MISHANMAUYRA Date: Fri, 13 Jun 2025 15:19:58 +0530 Subject: [PATCH 04/12] hipblas_matmul_fix --- csrc/ops.hip | 148 +++++++++++++++++++++++------------------------ csrc/ops_hip.cuh | 2 +- 2 files changed, 75 insertions(+), 75 deletions(-) diff --git a/csrc/ops.hip b/csrc/ops.hip index 93b27ed8e..60cf836b9 100644 --- a/csrc/ops.hip +++ b/csrc/ops.hip @@ -476,50 +476,50 @@ template int igemmlt( hipblasLtPointerMode_t pointerMode = HIPBLASLT_POINTER_MODE_ALPHA_DEVICE_VECTOR_BETA_HOST; has_error |= checkHipblasStatus(hipblasLtMatrixLayoutCreate(&aDesc, HIP_R_8I, m, k, lda)); - if(has_error != 0) - { - std::cout<<"failed to run hipblasLtMatrixLayoutCreate for Adesc:"< int igemmlt( fprintf(stderr, "Error: Matmul Algo Heuristic didn't return algorithms\n"); } else { int alpha = 1, beta = 0; - void* d_workspace = nullptr; - uint64_t workspace_size = 0; - for(int i = 0; i < returnedAlgoCount; i++) + void* d_workspace = nullptr; + uint64_t workspace_size = 0; + for(int i = 0; i < returnedAlgoCount; i++) workspace_size = max(workspace_size, heuristicResult[i].workspaceSize); - hipMalloc(&d_workspace, workspace_size); + hipMalloc(&d_workspace, workspace_size); has_error |= checkHipblasStatus(hipblasLtMatmul( ltHandle, matmulDesc, &alpha, A, aDesc, B, bDesc, &beta, (int32_t*)C, cDesc, (int32_t*)C, cDesc, - &heuristicResult[0].algo, d_workspace, workspace_size, stream + &heuristicResult[0].algo, d_workspace, workspace_size, stream )); hipFree(d_workspace); - if(has_error != 0) - { + if(has_error != 0) + { std::cout<<"failed to run hipblasLtMatmul"< int igemmlt( &pointerMode, sizeof(alphaVec) )); - if(has_error != 0) - { - std::cout<<"failed to run hipblasLtMatmulDescSetAttribute HIPBLASLT_MATMUL_DESC_POINTER_MODE for int8"< Date: Fri, 13 Jun 2025 15:33:03 +0530 Subject: [PATCH 05/12] Update ops.hip --- csrc/ops.hip | 12 ++---------- 1 file changed, 2 insertions(+), 10 deletions(-) diff --git a/csrc/ops.hip b/csrc/ops.hip index 60cf836b9..92c3ec273 100644 --- a/csrc/ops.hip +++ b/csrc/ops.hip @@ -583,18 +583,14 @@ template int igemmlt( if (!SCALE_ROWS) { float alpha = 1.0f, beta = 0.0f; - uint64_t workspace_size = 0; - void* d_workspace = nullptr; - hipMalloc(&d_workspace, workspace_size); has_error |= checkHipblasStatus(hipblasLtMatmul( ltHandle, matmulDesc, &alpha, A, aDesc, B, bDesc, &beta, (int8_t*)C, cDesc, (int8_t*)C, cDesc, - NULL, d_workspace, workspace_size, stream + NULL, NULL, 0, stream )); - hipFree(d_workspace); if(has_error != 0) { std::cout<<"failed to run hipblasLtMatmul with int8"< int igemmlt( return has_error; } - uint64_t workspace_size = 0; - void* d_workspace = nullptr; - hipMalloc(&d_workspace, workspace_size); has_error |= checkHipblasStatus(hipblasLtMatmul( ltHandle, matmulDesc, row_scale, A, aDesc, B, bDesc, &beta, (int8_t*)C, cDesc, (int8_t*)C, cDesc, - NULL, d_workspace, workspace_size, stream + NULL, NULL, 0, stream )); - hipFree(d_workspace); if(has_error != 0) { std::cout<<"failed to run hipblasLtMatmul with int8"< Date: Fri, 13 Jun 2025 15:52:22 +0530 Subject: [PATCH 06/12] Revert "Update ops.hip" This reverts commit 0a9a4915b37bea65a1556492840374e74f1219a7. --- csrc/ops.hip | 12 ++++++++++-- 1 file changed, 10 insertions(+), 2 deletions(-) diff --git a/csrc/ops.hip b/csrc/ops.hip index 92c3ec273..60cf836b9 100644 --- a/csrc/ops.hip +++ b/csrc/ops.hip @@ -583,14 +583,18 @@ template int igemmlt( if (!SCALE_ROWS) { float alpha = 1.0f, beta = 0.0f; + uint64_t workspace_size = 0; + void* d_workspace = nullptr; + hipMalloc(&d_workspace, workspace_size); has_error |= checkHipblasStatus(hipblasLtMatmul( ltHandle, matmulDesc, &alpha, A, aDesc, B, bDesc, &beta, (int8_t*)C, cDesc, (int8_t*)C, cDesc, - NULL, NULL, 0, stream + NULL, d_workspace, workspace_size, stream )); + hipFree(d_workspace); if(has_error != 0) { std::cout<<"failed to run hipblasLtMatmul with int8"< int igemmlt( return has_error; } + uint64_t workspace_size = 0; + void* d_workspace = nullptr; + hipMalloc(&d_workspace, workspace_size); has_error |= checkHipblasStatus(hipblasLtMatmul( ltHandle, matmulDesc, row_scale, A, aDesc, B, bDesc, &beta, (int8_t*)C, cDesc, (int8_t*)C, cDesc, - NULL, NULL, 0, stream + NULL, d_workspace, workspace_size, stream )); + hipFree(d_workspace); if(has_error != 0) { std::cout<<"failed to run hipblasLtMatmul with int8"< Date: Fri, 13 Jun 2025 20:40:33 +0530 Subject: [PATCH 07/12] Update ops_hip.cuh --- csrc/ops_hip.cuh | 15 --------------- 1 file changed, 15 deletions(-) diff --git a/csrc/ops_hip.cuh b/csrc/ops_hip.cuh index d803362fe..d6a221ecf 100644 --- a/csrc/ops_hip.cuh +++ b/csrc/ops_hip.cuh @@ -113,21 +113,6 @@ class Context }; -/* -class ContextLt -{ - public: - hipblasLtHandle_t m_handle; - - ContextLt() - { - hipblasLtHandle_t handle; - hipblasLtCreate(&handle); - m_handle = handle; - } -}; -*/ - class ContextHipsparse { public: From cd3eea7280a72e523923d1dd36b0dc168ff4cf97 Mon Sep 17 00:00:00 2001 From: MISHANMAURYA <118961433+MISHANMAURYA@users.noreply.github.com> Date: Fri, 13 Jun 2025 21:15:54 +0530 Subject: [PATCH 08/12] Update ops.hip --- csrc/ops.hip | 24 +++++------------------- 1 file changed, 5 insertions(+), 19 deletions(-) diff --git a/csrc/ops.hip b/csrc/ops.hip index 60cf836b9..831486624 100644 --- a/csrc/ops.hip +++ b/csrc/ops.hip @@ -485,7 +485,7 @@ template int igemmlt( has_error |= checkHipblasStatus(hipblasLtMatrixLayoutCreate(&bDesc, HIP_R_8I, m, n, ldb)); if(has_error != 0) { - std::cout<<"failed to run hipblasLtMatrixLayoutCreate for Bdesc:"< int igemmlt( has_error |= checkHipblasStatus(hipblasLtMatrixLayoutCreate(&cDesc, outType, k, n, ldc)); if(has_error != 0) { - std::cout<<"failed to run hipblasLtMatrixLayoutCreate for Cdesc"< int igemmlt( fprintf(stderr, "Error: Matmul Algo Heuristic didn't return algorithms\n"); } else { int alpha = 1, beta = 0; - void* d_workspace = nullptr; - uint64_t workspace_size = 0; - for(int i = 0; i < returnedAlgoCount; i++) - workspace_size = max(workspace_size, heuristicResult[i].workspaceSize); - hipMalloc(&d_workspace, workspace_size); has_error |= checkHipblasStatus(hipblasLtMatmul( ltHandle, matmulDesc, &alpha, A, aDesc, B, bDesc, &beta, (int32_t*)C, cDesc, (int32_t*)C, cDesc, - &heuristicResult[0].algo, d_workspace, workspace_size, stream + &heuristicResult[0].algo, NULL, 0, stream )); - hipFree(d_workspace); if(has_error != 0) { std::cout<<"failed to run hipblasLtMatmul"< int igemmlt( if (!SCALE_ROWS) { float alpha = 1.0f, beta = 0.0f; - uint64_t workspace_size = 0; - void* d_workspace = nullptr; - hipMalloc(&d_workspace, workspace_size); has_error |= checkHipblasStatus(hipblasLtMatmul( ltHandle, matmulDesc, &alpha, A, aDesc, B, bDesc, &beta, (int8_t*)C, cDesc, (int8_t*)C, cDesc, - NULL, d_workspace, workspace_size, stream + NULL, NULL, 0, stream )); - hipFree(d_workspace); if(has_error != 0) { std::cout<<"failed to run hipblasLtMatmul with int8"< int igemmlt( return has_error; } - uint64_t workspace_size = 0; - void* d_workspace = nullptr; - hipMalloc(&d_workspace, workspace_size); has_error |= checkHipblasStatus(hipblasLtMatmul( ltHandle, matmulDesc, row_scale, A, aDesc, B, bDesc, &beta, (int8_t*)C, cDesc, (int8_t*)C, cDesc, - NULL, d_workspace, workspace_size, stream + NULL, NULL, 0, stream )); - hipFree(d_workspace); if(has_error != 0) { std::cout<<"failed to run hipblasLtMatmul with int8"< Date: Fri, 13 Jun 2025 21:19:10 +0530 Subject: [PATCH 09/12] Lint --- csrc/ops.hip | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/csrc/ops.hip b/csrc/ops.hip index 831486624..91b17cdd8 100644 --- a/csrc/ops.hip +++ b/csrc/ops.hip @@ -560,7 +560,7 @@ template int igemmlt( B, bDesc, &beta, (int32_t*)C, cDesc, (int32_t*)C, cDesc, - &heuristicResult[0].algo, NULL, 0, stream + &heuristicResult[0].algo, NULL, 0, stream )); if(has_error != 0) { @@ -619,7 +619,7 @@ template int igemmlt( B, bDesc, &beta, (int8_t*)C, cDesc, (int8_t*)C, cDesc, - NULL, NULL, 0, stream + NULL, NULL, 0, stream )); if(has_error != 0) { From 49abe7e440c2ee6f82bf2c157d320b7bbbb3449a Mon Sep 17 00:00:00 2001 From: MISHANMAURYA <118961433+MISHANMAURYA@users.noreply.github.com> Date: Fri, 13 Jun 2025 21:40:46 +0530 Subject: [PATCH 10/12] Update ops_hip.cuh --- csrc/ops_hip.cuh | 15 ++++++++------- 1 file changed, 8 insertions(+), 7 deletions(-) diff --git a/csrc/ops_hip.cuh b/csrc/ops_hip.cuh index d6a221ecf..80ef86ba5 100644 --- a/csrc/ops_hip.cuh +++ b/csrc/ops_hip.cuh @@ -113,19 +113,20 @@ class Context }; -class ContextHipsparse +/* +class ContextLt { public: - hipsparseHandle_t m_handle; - - ContextHipsparse() + hipblasLtHandle_t m_handle; + ContextLt() { - hipsparseHandle_t handle; - hipsparseCreate(&handle); + hipblasLtHandle_t handle; + hipblasLtCreate(&handle); m_handle = handle; } - }; +*/ + template void estimateQuantiles(T *A, float *code, float offset, int n); From d547a10852d80230aef2113b3747e2758c12b9d3 Mon Sep 17 00:00:00 2001 From: MISHANMAURYA <118961433+MISHANMAURYA@users.noreply.github.com> Date: Fri, 13 Jun 2025 21:51:22 +0530 Subject: [PATCH 11/12] Update ops_hip.cuh --- csrc/ops_hip.cuh | 15 +++++++-------- 1 file changed, 7 insertions(+), 8 deletions(-) diff --git a/csrc/ops_hip.cuh b/csrc/ops_hip.cuh index 80ef86ba5..d6a221ecf 100644 --- a/csrc/ops_hip.cuh +++ b/csrc/ops_hip.cuh @@ -113,20 +113,19 @@ class Context }; -/* -class ContextLt +class ContextHipsparse { public: - hipblasLtHandle_t m_handle; - ContextLt() + hipsparseHandle_t m_handle; + + ContextHipsparse() { - hipblasLtHandle_t handle; - hipblasLtCreate(&handle); + hipsparseHandle_t handle; + hipsparseCreate(&handle); m_handle = handle; } -}; -*/ +}; template void estimateQuantiles(T *A, float *code, float offset, int n); From da951d820531fbef492be95fcba839e34e11bf11 Mon Sep 17 00:00:00 2001 From: MISHANMAURYA <118961433+MISHANMAURYA@users.noreply.github.com> Date: Fri, 13 Jun 2025 22:09:21 +0530 Subject: [PATCH 12/12] Update ops_hip.cuh --- csrc/ops_hip.cuh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/csrc/ops_hip.cuh b/csrc/ops_hip.cuh index d6a221ecf..b1136dda9 100644 --- a/csrc/ops_hip.cuh +++ b/csrc/ops_hip.cuh @@ -99,7 +99,7 @@ typedef enum Funcs_t class Context { public: - hipblasLtHandle_t m_handle; + hipblasLtHandle_t m_handle; //rocblas_handle m_handle; Context()