Skip to content

Qinglei/trmm gpu #154

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
QingleiCao merged 2 commits into from
Mar 11, 2026
Merged

Qinglei/trmm gpu #154

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
27a4fd2
Update PaRSEC
May 22, 2025
b2589e5
Add GPU support in trmm
QingleiCao Feb 13, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 1 addition & 1 deletion parsec
Open in desktop
Submodule parsec updated 39 files
+0 −5 .github/CI/build.sh
+0 −7 .github/CI/configure.sh
+16 −11 .github/CI/env_setup.sh
+0 −5 .github/CI/install.sh
+0 −5 .github/CI/setup.sh
+0 −97 .github/CI/spack_setup.sh
+0 −14 .github/CI/test.sh
+0 −213 .github/workflows/build_cmake.yml
+35 −12 .github/workflows/main.yml
+7 −7 CMakeLists.txt
+1 −0 cmake_modules/PaRSECConfig.cmake.in
+4 −0 configure
+9 −1 parsec/class/list.h
+15 −0 parsec/class/parsec_object.c
+19 −8 parsec/class/parsec_object.h
+6 −0 parsec/data.c
+4 −0 parsec/data.h
+3 −2 parsec/data_dist/matrix/CMakeLists.txt
+5 −5 parsec/data_dist/matrix/apply_wrapper.c
+3 −1 parsec/data_internal.h
+1 −1 parsec/include/parsec/sys/atomic-x86_64.h
+35 −44 parsec/interfaces/ptg/ptg-compiler/jdf2c.c
+30 −10 parsec/interfaces/ptg/ptg-compiler/parsec.y
+30 −26 parsec/mca/device/cuda/device_cuda_module.c
+75 −9 parsec/mca/device/device_gpu.c
+5 −1 parsec/mca/device/hip/ValidateModule.CMake
+5 −0 parsec/mca/device/transfer_gpu.c
+19 −4 parsec/parsec.c
+8 −11 parsec/parsec_mpi_funnelled.c
+2 −1 parsec/runtime.h
+34 −19 parsec/scheduling.c
+17 −0 parsec/scheduling.h
+1 −1 parsec/utils/argv.c
+1 −1 parsec/utils/process_name.c
+22 −15 parsec/utils/zone_malloc.c
+3 −1 parsec/utils/zone_malloc.h
+1 −1 tests/dsl/ptg/recursive.jdf
+6 −0 tests/runtime/cuda/nvlink_main.c
+8 −4 tools/profiling/python/CMakeLists.txt
62 changes: 62 additions & 0 deletions src/ztrmm_LLN.jdf
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -8,6 +8,10 @@ extern "C" %{
* @precisions normal z -> s d c
*
*/
#include "dplasma/config.h"
#if defined(DPLASMA_HAVE_CUDA)
#include <cublas.h>
#endif /* defined(DPLASMA_HAVE_CUDA) */
#include "dplasmajdf.h"
#include "parsec/data_dist/matrix/matrix.h"

Expand Down Expand Up @@ -54,6 +58,9 @@ descA [type = "const parsec_tiled_matrix_t*" hidden = on default = "((dplas
ddescB [type = "dplasma_data_collection_t*"]
descB [type = "parsec_tiled_matrix_t*" hidden = on default = "((dplasma_data_collection_t*)ddescB)->dc_original" aligned=ddescB]

hip_handles_infokey [type = "int" hidden = on default = -1 ]


read_A(m, k) [profile = off]
/* Execution Space */
m = 0..(descB->mt-1)
Expand Down Expand Up @@ -153,6 +160,61 @@ loc_C = %{ return LOC(descB, (descB->mt-1)-m, n); %}
type_data = %{ return ADTT_DC(ddescB, loc_C, B_SHAPE, LAPACK); %} ]
-> ((k+m) < (descB->mt-2)) ? C zgemm(m, n, k+1) /* dep OUT: rely on datacopy dtt for sending */


BODY [type=CUDA]
{
#if defined(PRECISION_z) || defined(PRECISION_c)
cuDoubleComplex lalpha = make_cuDoubleComplex(creal(alpha), cimag(alpha));
cuDoubleComplex lbeta = make_cuDoubleComplex( 1., 0.);
#else
double lalpha = alpha;
double lbeta = 1.0;
#endif
int tempmm = (((descB->mt-1)-m)==(descB->mt-1)) ? (descB->m-(((descB->mt-1)-m)*descB->mb)) : descB->mb;
int tempnn = ((n)==(descB->nt-1)) ? (descB->n-(n*descB->nb)) : descB->nb;
int lda = LDA(ddescA, A);
int ldb = LDA(ddescB, B);
int ldc = LDA(ddescB, C);

cublasStatus_t status;
cublasSetKernelStream( parsec_body.stream );
cublasZgemm( dplasma_lapack_const(trans), 'N',
tempmm, tempnn, descB->mb,
lalpha, (cuDoubleComplex*)A, lda,
(cuDoubleComplex*)B, ldb,
lbeta, (cuDoubleComplex*)C, ldc );
status = cublasGetError();
PARSEC_CUDA_CHECK_ERROR( "cublasZgemm ", status, {return PARSEC_HOOK_RETURN_ERROR;} );
}
END

BODY [type=HIP]
{
#if defined(PRECISION_z) || defined(PRECISION_c)
hipDoubleComplex lalpha = make_hipDoubleComplex(creal(alpha), cimag(alpha));
hipDoubleComplex lbeta = make_hipDoubleComplex( 1., 0.);
#else
double lalpha = alpha;
double lbeta = 1.0;
#endif
int tempmm = (((descB->mt-1)-m)==(descB->mt-1)) ? (descB->m-(((descB->mt-1)-m)*descB->mb)) : descB->mb;
int tempnn = ((n)==(descB->nt-1)) ? (descB->n-(n*descB->nb)) : descB->nb;
int lda = LDA(ddescA, A);
int ldb = LDA(ddescB, B);
int ldc = LDA(ddescB, C);

hipblasStatus_t status;
dplasma_hip_handles_t *handles = parsec_info_get(&gpu_stream->infos, hip_handles_infokey);
assert(NULL != handles);
status = hipblasZgemm( handles->hipblas_handle, dplasma_hipblas_op(trans), HIPBLAS_OP_N,
tempmm, tempnn, descB->mb,
&lalpha, (hipDoubleComplex*)A, lda,
(hipDoubleComplex*)B, ldb,
&lbeta, (hipDoubleComplex*)C, ldc );
DPLASMA_HIPBLAS_CHECK_ERROR( "hipblasZgemm ", status, {return PARSEC_HOOK_RETURN_ERROR;} );
}
END

BODY
{
int tempmm = (((descB->mt-1)-m)==(descB->mt-1)) ? (descB->m-(((descB->mt-1)-m)*descB->mb)) : descB->mb;
Expand Down
62 changes: 62 additions & 0 deletions src/ztrmm_LLT.jdf
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -8,6 +8,10 @@ extern "C" %{
* @precisions normal z -> s d c
*
*/
#include "dplasma/config.h"
#if defined(DPLASMA_HAVE_CUDA)
#include <cublas.h>
#endif /* defined(DPLASMA_HAVE_CUDA) */
#include "dplasmajdf.h"
#include "parsec/data_dist/matrix/matrix.h"

Expand Down Expand Up @@ -54,6 +58,8 @@ descA [type = "const parsec_tiled_matrix_t*" hidden = on default = "((dplas
ddescB [type = "dplasma_data_collection_t*"]
descB [type = "parsec_tiled_matrix_t*" hidden = on default = "((dplasma_data_collection_t*)ddescB)->dc_original" aligned=ddescB]

hip_handles_infokey [type = "int" hidden = on default = -1 ]

read_A(m, k) [profile = off]
/* Execution Space */
m = 0 .. (descB->mt-1)
Expand Down Expand Up @@ -154,6 +160,62 @@ loc_C = %{ return LOC(descB, m, n); %}
type_data = %{ return ADTT_DC(ddescB, loc_C, B_SHAPE, LAPACK); %} ]
-> (k < (descB->mt-1)) ? C zgemm(m, n, k+1) /* dep OUT: rely on datacopy dtt for sending */

BODY [type=CUDA]
{
#if defined(PRECISION_z) || defined(PRECISION_c)
cuDoubleComplex lalpha = make_cuDoubleComplex(creal(alpha), cimag(alpha));
cuDoubleComplex lbeta = make_cuDoubleComplex( 1., 0.);
#else
double lalpha = alpha;
double lbeta = 1.0;
#endif
int tempmm = ((m)==(descB->mt-1)) ? (descB->m-(m*descB->mb)) : descB->mb;
int tempnn = ((n)==(descB->nt-1)) ? (descB->n-(n*descB->nb)) : descB->nb;
int tempkm = ((k)==(descA->mt-1)) ? (descA->m-(k*descA->mb)) : descA->mb;
int lda = LDA(ddescA, A);
int ldb = LDA(ddescB, B);
int ldc = LDA(ddescB, C);

cublasStatus_t status;
cublasSetKernelStream( parsec_body.stream );
cublasZgemm( dplasma_lapack_const(trans), 'N',
tempmm, tempnn, tempkm,
lalpha, (cuDoubleComplex*)A, lda,
(cuDoubleComplex*)B, ldb,
lbeta, (cuDoubleComplex*)C, ldc );
status = cublasGetError();
PARSEC_CUDA_CHECK_ERROR( "cublasZgemm ", status, {return PARSEC_HOOK_RETURN_ERROR;} );
}
END

BODY [type=HIP]
{
#if defined(PRECISION_z) || defined(PRECISION_c)
hipDoubleComplex lalpha = make_hipDoubleComplex(creal(alpha), cimag(alpha));
hipDoubleComplex lbeta = make_hipDoubleComplex( 1., 0.);
#else
double lalpha = alpha;
double lbeta = 1.0;
#endif
int tempmm = ((m)==(descB->mt-1)) ? (descB->m-(m*descB->mb)) : descB->mb;
int tempnn = ((n)==(descB->nt-1)) ? (descB->n-(n*descB->nb)) : descB->nb;
int tempkm = ((k)==(descA->mt-1)) ? (descA->m-(k*descA->mb)) : descA->mb;
int lda = LDA(ddescA, A);
int ldb = LDA(ddescB, B);
int ldc = LDA(ddescB, C);

hipblasStatus_t status;
dplasma_hip_handles_t *handles = parsec_info_get(&gpu_stream->infos, hip_handles_infokey);
assert(NULL != handles);
status = hipblasZgemm( handles->hipblas_handle, dplasma_hipblas_op(trans), HIPBLAS_OP_N,
tempmm, tempnn, tempkm,
&lalpha, (hipDoubleComplex*)A, lda,
(hipDoubleComplex*)B, ldb,
&lbeta, (hipDoubleComplex*)C, ldc );
DPLASMA_HIPBLAS_CHECK_ERROR( "hipblasZgemm ", status, {return PARSEC_HOOK_RETURN_ERROR;} );
}
END

BODY
{
int tempmm = ((m)==(descB->mt-1)) ? (descB->m-(m*descB->mb)) : descB->mb;
Expand Down
62 changes: 62 additions & 0 deletions src/ztrmm_LUN.jdf
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -8,6 +8,10 @@ extern "C" %{
* @precisions normal z -> s d c
*
*/
#include "dplasma/config.h"
#if defined(DPLASMA_HAVE_CUDA)
#include <cublas.h>
#endif /* defined(DPLASMA_HAVE_CUDA) */
#include "dplasmajdf.h"
#include "parsec/data_dist/matrix/matrix.h"

Expand Down Expand Up @@ -54,6 +58,8 @@ descA [type = "const parsec_tiled_matrix_t*" hidden = on default = "((dplas
ddescB [type = "dplasma_data_collection_t*"]
descB [type = "parsec_tiled_matrix_t*" hidden = on default = "((dplasma_data_collection_t*)ddescB)->dc_original" aligned=ddescB]

hip_handles_infokey [type = "int" hidden = on default = -1 ]

read_A(m, k) [profile = off]
/* Execution Space */
m = 0 .. (descB->mt-1)
Expand Down Expand Up @@ -154,6 +160,62 @@ loc_C = %{ return LOC(descB, m, n); %}
type_data = %{ return ADTT_DC(ddescB, loc_C, B_SHAPE, LAPACK); %} ]
-> (k < (descB->mt-1)) ? C zgemm(m, n, k+1) /* dep OUT: rely on datacopy dtt for sending */

BODY [type=CUDA]
{
#if defined(PRECISION_z) || defined(PRECISION_c)
cuDoubleComplex lalpha = make_cuDoubleComplex(creal(alpha), cimag(alpha));
cuDoubleComplex lbeta = make_cuDoubleComplex( 1., 0.);
#else
double lalpha = alpha;
double lbeta = 1.0;
#endif
int tempmm = ((m)==(descB->mt-1)) ? (descB->m-(m*descB->mb)) : descB->mb;
int tempnn = ((n)==(descB->nt-1)) ? (descB->n-(n*descB->nb)) : descB->nb;
int tempkn = ((k)==(descA->nt-1)) ? (descA->n-(k*descA->nb)) : descA->nb;
int lda = LDA(ddescA, A);
int ldb = LDA(ddescB, B);
int ldc = LDA(ddescB, C);

cublasStatus_t status;
cublasSetKernelStream( parsec_body.stream );
cublasZgemm( dplasma_lapack_const(trans), 'N',
tempmm, tempnn, tempkn,
lalpha, (cuDoubleComplex*)A, lda,
(cuDoubleComplex*)B, ldb,
lbeta, (cuDoubleComplex*)C, ldc );
status = cublasGetError();
PARSEC_CUDA_CHECK_ERROR( "cublasZgemm ", status, {return PARSEC_HOOK_RETURN_ERROR;} );
}
END

BODY [type=HIP]
{
#if defined(PRECISION_z) || defined(PRECISION_c)
hipDoubleComplex lalpha = make_hipDoubleComplex(creal(alpha), cimag(alpha));
hipDoubleComplex lbeta = make_hipDoubleComplex( 1., 0.);
#else
double lalpha = alpha;
double lbeta = 1.0;
#endif
int tempmm = ((m)==(descB->mt-1)) ? (descB->m-(m*descB->mb)) : descB->mb;
int tempnn = ((n)==(descB->nt-1)) ? (descB->n-(n*descB->nb)) : descB->nb;
int tempkn = ((k)==(descA->nt-1)) ? (descA->n-(k*descA->nb)) : descA->nb;
int lda = LDA(ddescA, A);
int ldb = LDA(ddescB, B);
int ldc = LDA(ddescB, C);

hipblasStatus_t status;
dplasma_hip_handles_t *handles = parsec_info_get(&gpu_stream->infos, hip_handles_infokey);
assert(NULL != handles);
status = hipblasZgemm( handles->hipblas_handle, dplasma_hipblas_op(trans), HIPBLAS_OP_N,
tempmm, tempnn, tempkn,
&lalpha, (hipDoubleComplex*)A, lda,
(hipDoubleComplex*)B, ldb,
&lbeta, (hipDoubleComplex*)C, ldc );
DPLASMA_HIPBLAS_CHECK_ERROR( "hipblasZgemm ", status, {return PARSEC_HOOK_RETURN_ERROR;} );
}
END

BODY
{
int tempmm = ((m)==(descB->mt-1)) ? (descB->m-(m*descB->mb)) : descB->mb;
Expand Down
60 changes: 60 additions & 0 deletions src/ztrmm_LUT.jdf
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -8,6 +8,10 @@ extern "C" %{
* @precisions normal z -> s d c
*
*/
#include "dplasma/config.h"
#if defined(DPLASMA_HAVE_CUDA)
#include <cublas.h>
#endif /* defined(DPLASMA_HAVE_CUDA) */
#include "dplasmajdf.h"
#include "parsec/data_dist/matrix/matrix.h"

Expand Down Expand Up @@ -54,6 +58,8 @@ descA [type = "const parsec_tiled_matrix_t*" hidden = on default = "((dplas
ddescB [type = "dplasma_data_collection_t*"]
descB [type = "parsec_tiled_matrix_t*" hidden = on default = "((dplasma_data_collection_t*)ddescB)->dc_original" aligned=ddescB]

hip_handles_infokey [type = "int" hidden = on default = -1 ]

read_A(m, k) [profile = off]
/* Execution Space */
m = 0..(descB->mt-1)
Expand Down Expand Up @@ -153,6 +159,60 @@ loc_C = %{ return LOC(descB, (descB->mt-1)-m, n); %}
type_data = %{ return ADTT_DC(ddescB, loc_C, B_SHAPE, LAPACK); %} ]
-> ((k+m) < (descB->mt-2)) ? C zgemm(m, n, k+1) /* dep OUT: rely on datacopy dtt for sending */

BODY [type=CUDA]
{
#if defined(PRECISION_z) || defined(PRECISION_c)
cuDoubleComplex lalpha = make_cuDoubleComplex(creal(alpha), cimag(alpha));
cuDoubleComplex lbeta = make_cuDoubleComplex( 1., 0.);
#else
double lalpha = alpha;
double lbeta = 1.0;
#endif
int tempmm = (((descB->mt-1)-m)==(descB->mt-1)) ? (descB->m-(((descB->mt-1)-m)*descB->mb)) : descB->mb;
int tempnn = ((n)==(descB->nt-1)) ? (descB->n-(n*descB->nb)) : descB->nb;
int lda = LDA(ddescA, A);
int ldb = LDA(ddescB, B);
int ldc = LDA(ddescB, C);

cublasStatus_t status;
cublasSetKernelStream( parsec_body.stream );
cublasZgemm( dplasma_lapack_const(trans), 'N',
tempmm, tempnn, descB->mb,
lalpha, (cuDoubleComplex*)A, lda,
(cuDoubleComplex*)B, ldb,
lbeta, (cuDoubleComplex*)C, ldc );
status = cublasGetError();
PARSEC_CUDA_CHECK_ERROR( "cublasZgemm ", status, {return PARSEC_HOOK_RETURN_ERROR;} );
}
END

BODY [type=HIP]
{
#if defined(PRECISION_z) || defined(PRECISION_c)
hipDoubleComplex lalpha = make_hipDoubleComplex(creal(alpha), cimag(alpha));
hipDoubleComplex lbeta = make_hipDoubleComplex( 1., 0.);
#else
double lalpha = alpha;
double lbeta = 1.0;
#endif
int tempmm = (((descB->mt-1)-m)==(descB->mt-1)) ? (descB->m-(((descB->mt-1)-m)*descB->mb)) : descB->mb;
int tempnn = ((n)==(descB->nt-1)) ? (descB->n-(n*descB->nb)) : descB->nb;
int lda = LDA(ddescA, A);
int ldb = LDA(ddescB, B);
int ldc = LDA(ddescB, C);

hipblasStatus_t status;
dplasma_hip_handles_t *handles = parsec_info_get(&gpu_stream->infos, hip_handles_infokey);
assert(NULL != handles);
status = hipblasZgemm( handles->hipblas_handle, dplasma_hipblas_op(trans), HIPBLAS_OP_N,
tempmm, tempnn, descB->mb,
&lalpha, (hipDoubleComplex*)A, lda,
(hipDoubleComplex*)B, ldb,
&lbeta, (hipDoubleComplex*)C, ldc );
DPLASMA_HIPBLAS_CHECK_ERROR( "hipblasZgemm ", status, {return PARSEC_HOOK_RETURN_ERROR;} );
}
END

BODY
{
int tempmm = (((descB->mt-1)-m)==(descB->mt-1)) ? (descB->m-(((descB->mt-1)-m)*descB->mb)) : descB->mb;
Expand Down
Loading
Loading