Skip to content

Removed unnecessary branching in dot products #469

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

Closed
shakedregev wants to merge 2 commits into from
Closed

Removed unnecessary branching in dot products #469

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
ed85271
fixed some dot product issues
Aug 11, 2026
12ad9e1
Apply pre-commmit fixes
shakedregev Aug 11, 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
84 changes: 30 additions & 54 deletions resolve/vector/VectorHandlerCuda.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -235,32 +235,21 @@ namespace ReSolve
vector::Vector* y)
{
using namespace constants;
if (k < 200)
{
cuda::axpy_multi(size,
k,
x->getData(memory::DEVICE),
y->getData(memory::DEVICE),
alpha->getData(memory::DEVICE));
}
else
{
cublasHandle_t handle_cublas = workspace_->getCublasHandle();
cublasDgemm(handle_cublas,
CUBLAS_OP_N,
CUBLAS_OP_N,
size, // m
1, // n
k, // k
&MINUS_ONE, // alpha
x->getData(memory::DEVICE), // A
size, // lda
alpha->getData(memory::DEVICE), // B
k, // ldb
&ONE,
y->getData(memory::DEVICE), // c
size); // ldc
}
cublasHandle_t handle_cublas = workspace_->getCublasHandle();
cublasDgemm(handle_cublas,
CUBLAS_OP_N,
CUBLAS_OP_N,
size, // m
1, // n
k, // k
&MINUS_ONE, // alpha
x->getData(memory::DEVICE), // A
size, // lda
alpha->getData(memory::DEVICE), // B
k, // ldb
&ONE,
y->getData(memory::DEVICE), // c
size); // ldc
y->setDataUpdated(memory::DEVICE);
}

Expand All @@ -285,34 +274,21 @@ namespace ReSolve
vector::Vector* res)
{
using namespace constants;

if (k < 200)
{
cuda::dot_2_multi(size,
k,
x->getData(0, memory::DEVICE),
x->getData(1, memory::DEVICE),
V->getData(memory::DEVICE),
res->getData(memory::DEVICE));
}
else
{
cublasHandle_t handle_cublas = workspace_->getCublasHandle();
cublasDgemm(handle_cublas,
CUBLAS_OP_T,
CUBLAS_OP_N,
k, // m
2, // n
size, // k
&ONE, // alpha
V->getData(memory::DEVICE), // A
size, // lda
x->getData(memory::DEVICE), // B
size, // ldb
&ZERO,
res->getData(memory::DEVICE), // c
k); // ldc
}
cublasHandle_t handle_cublas = workspace_->getCublasHandle();
cublasDgemm(handle_cublas,
CUBLAS_OP_T,
CUBLAS_OP_N,
k, // m
2, // n
size, // k
&ONE, // alpha
V->getData(memory::DEVICE), // A
size, // lda
x->getData(memory::DEVICE), // B
size, // ldb
&ZERO,
res->getData(memory::DEVICE), // c
k); // ldc
res->setDataUpdated(memory::DEVICE);
}

Expand Down
84 changes: 30 additions & 54 deletions resolve/vector/VectorHandlerHip.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -234,32 +234,21 @@ namespace ReSolve
vector::Vector* y)
{
using namespace constants;
if (k < 200)
{
hip::axpy_multi(size,
k,
x->getData(memory::DEVICE),
y->getData(memory::DEVICE),
alpha->getData(memory::DEVICE));
}
else
{
rocblas_handle handle_rocblas = workspace_->getRocblasHandle();
rocblas_dgemm(handle_rocblas,
rocblas_operation_none,
rocblas_operation_none,
size, // m
1, // n
k, // k
&MINUS_ONE, // alpha
x->getData(memory::DEVICE), // A
size, // lda
alpha->getData(memory::DEVICE), // B
k, // ldb
&ONE,
y->getData(memory::DEVICE), // c
size); // ldc
}
rocblas_handle handle_rocblas = workspace_->getRocblasHandle();
rocblas_dgemm(handle_rocblas,
rocblas_operation_none,
rocblas_operation_none,
size, // m
1, // n
k, // k
&MINUS_ONE, // alpha
x->getData(memory::DEVICE), // A
size, // lda
alpha->getData(memory::DEVICE), // B
k, // ldb
&ONE,
y->getData(memory::DEVICE), // c
size); // ldc
y->setDataUpdated(memory::DEVICE);
mem_.deviceSynchronize();
}
Expand All @@ -285,34 +274,21 @@ namespace ReSolve
vector::Vector* res)
{
using namespace constants;

if (k < 200)
{
hip::dot_2_multi(size,
k,
x->getData(0, memory::DEVICE),
x->getData(1, memory::DEVICE),
V->getData(memory::DEVICE),
res->getData(memory::DEVICE));
}
else
{
rocblas_handle handle_rocblas = workspace_->getRocblasHandle();
rocblas_dgemm(handle_rocblas,
rocblas_operation_transpose,
rocblas_operation_none,
k, // m
2, // n
size, // k
&ONE, // alpha
V->getData(memory::DEVICE), // A
size, // lda
x->getData(memory::DEVICE), // B
size, // ldb
&ZERO,
res->getData(memory::DEVICE), // c
k); // ldc
}
rocblas_handle handle_rocblas = workspace_->getRocblasHandle();
rocblas_dgemm(handle_rocblas,
rocblas_operation_transpose,
rocblas_operation_none,
k, // m
2, // n
size, // k
&ONE, // alpha
V->getData(memory::DEVICE), // A
size, // lda
x->getData(memory::DEVICE), // B
size, // ldb
&ZERO,
res->getData(memory::DEVICE), // c
k); // ldc
res->setDataUpdated(memory::DEVICE);
mem_.deviceSynchronize();
}
Expand Down