Skip to content

HIP Vec Filter #2016

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
jeremylt merged 2 commits into from
Aug 27, 2026
Merged

HIP Vec Filter #2016

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
19ecdd5
hip - mirror cuda vec filter
jeremylt Aug 27, 2026
6e0b742
gpu - mark gpu backend cpu functions with CeedPragmaSIMD
jeremylt Aug 27, 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
16 changes: 8 additions & 8 deletions backends/cuda-ref/ceed-cuda-ref-vector.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -228,7 +228,7 @@ static int CeedVectorSetArray_Cuda(const CeedVector vec, const CeedMemType mem_t
// Copy host array to value strided
//------------------------------------------------------------------------------
static int CeedHostCopyStrided_Cuda(CeedScalar *h_array, CeedSize start, CeedSize stop, CeedSize step, CeedScalar *h_copy_array) {
for (CeedSize i = start; i < stop; i += step) h_copy_array[i] = h_array[i];
CeedPragmaSIMD for (CeedSize i = start; i < stop; i += step) h_copy_array[i] = h_array[i];
return CEED_ERROR_SUCCESS;
}

Expand Down Expand Up @@ -292,7 +292,7 @@ static int CeedVectorCopyStrided_Cuda(CeedVector vec, CeedSize start, CeedSize s
// Set host array to value
//------------------------------------------------------------------------------
static int CeedHostSetValue_Cuda(CeedScalar *h_array, CeedSize length, CeedScalar val) {
for (CeedSize i = 0; i < length; i++) h_array[i] = val;
CeedPragmaSIMD for (CeedSize i = 0; i < length; i++) h_array[i] = val;
return CEED_ERROR_SUCCESS;
}

Expand Down Expand Up @@ -342,7 +342,7 @@ static int CeedVectorSetValue_Cuda(CeedVector vec, CeedScalar val) {
// Set host array to value strided
//------------------------------------------------------------------------------
static int CeedHostSetValueStrided_Cuda(CeedScalar *h_array, CeedSize start, CeedSize stop, CeedSize step, CeedScalar val) {
for (CeedSize i = start; i < stop; i += step) h_array[i] = val;
CeedPragmaSIMD for (CeedSize i = start; i < stop; i += step) h_array[i] = val;
return CEED_ERROR_SUCCESS;
}

Expand Down Expand Up @@ -647,7 +647,7 @@ static int CeedVectorNorm_Cuda(CeedVector vec, CeedNormType type, CeedScalar *no
// Take reciprocal of a vector on host
//------------------------------------------------------------------------------
static int CeedHostReciprocal_Cuda(CeedScalar *h_array, CeedSize length) {
for (CeedSize i = 0; i < length; i++) {
CeedPragmaSIMD for (CeedSize i = 0; i < length; i++) {
if (fabs(h_array[i]) > CEED_EPSILON) h_array[i] = 1. / h_array[i];
}
return CEED_ERROR_SUCCESS;
Expand Down Expand Up @@ -677,7 +677,7 @@ static int CeedVectorReciprocal_Cuda(CeedVector vec) {
// Compute x = alpha x on the host
//------------------------------------------------------------------------------
static int CeedHostScale_Cuda(CeedScalar *x_array, CeedScalar alpha, CeedSize length) {
for (CeedSize i = 0; i < length; i++) x_array[i] *= alpha;
CeedPragmaSIMD for (CeedSize i = 0; i < length; i++) x_array[i] *= alpha;
return CEED_ERROR_SUCCESS;
}

Expand Down Expand Up @@ -751,7 +751,7 @@ static int CeedVectorFilter_Cuda(CeedVector vec, CeedScalar threshold) {
// Compute y = alpha x + y on the host
//------------------------------------------------------------------------------
static int CeedHostAXPY_Cuda(CeedScalar *y_array, CeedScalar alpha, CeedScalar *x_array, CeedSize length) {
for (CeedSize i = 0; i < length; i++) y_array[i] += alpha * x_array[i];
CeedPragmaSIMD for (CeedSize i = 0; i < length; i++) y_array[i] += alpha * x_array[i];
return CEED_ERROR_SUCCESS;
}

Expand Down Expand Up @@ -798,7 +798,7 @@ static int CeedVectorAXPY_Cuda(CeedVector y, CeedScalar alpha, CeedVector x) {
// Compute y = alpha x + beta y on the host
//------------------------------------------------------------------------------
static int CeedHostAXPBY_Cuda(CeedScalar *y_array, CeedScalar alpha, CeedScalar beta, CeedScalar *x_array, CeedSize length) {
for (CeedSize i = 0; i < length; i++) y_array[i] = alpha * x_array[i] + beta * y_array[i];
CeedPragmaSIMD for (CeedSize i = 0; i < length; i++) y_array[i] = alpha * x_array[i] + beta * y_array[i];
return CEED_ERROR_SUCCESS;
}

Expand Down Expand Up @@ -833,7 +833,7 @@ static int CeedVectorAXPBY_Cuda(CeedVector y, CeedScalar alpha, CeedScalar beta,
// Compute the pointwise multiplication w = x .* y on the host
//------------------------------------------------------------------------------
static int CeedHostPointwiseMult_Cuda(CeedScalar *w_array, CeedScalar *x_array, CeedScalar *y_array, CeedSize length) {
for (CeedSize i = 0; i < length; i++) w_array[i] = x_array[i] * y_array[i];
CeedPragmaSIMD for (CeedSize i = 0; i < length; i++) w_array[i] = x_array[i] * y_array[i];
return CEED_ERROR_SUCCESS;
}

Expand Down
45 changes: 38 additions & 7 deletions backends/hip-ref/ceed-hip-ref-vector.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -281,7 +281,7 @@ static int CeedVectorSetArray_Hip(const CeedVector vec, const CeedMemType mem_ty
// Copy host array to value strided
//------------------------------------------------------------------------------
static int CeedHostCopyStrided_Hip(CeedScalar *h_array, CeedSize start, CeedSize stop, CeedSize step, CeedScalar *h_copy_array) {

@zatkins-dev zatkins-dev Aug 27, 2026

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Are these host functions actually tested? I think the way we set up the tests, they may not be. We should probably double-check.

@jeremylt jeremylt Aug 27, 2026

Copy link
Copy Markdown
Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

That would probably be a good student issue - modify the tests suite to ensure all these Host functions are tested

for (CeedSize i = start; i < stop; i += step) h_copy_array[i] = h_array[i];
CeedPragmaSIMD for (CeedSize i = start; i < stop; i += step) h_copy_array[i] = h_array[i];
return CEED_ERROR_SUCCESS;
}

Expand Down Expand Up @@ -347,7 +347,7 @@ static int CeedVectorCopyStrided_Hip(CeedVector vec, CeedSize start, CeedSize st
// Set host array to value
//------------------------------------------------------------------------------
static int CeedHostSetValue_Hip(CeedScalar *h_array, CeedSize length, CeedScalar val) {
for (CeedSize i = 0; i < length; i++) h_array[i] = val;
CeedPragmaSIMD for (CeedSize i = 0; i < length; i++) h_array[i] = val;
return CEED_ERROR_SUCCESS;
}

Expand Down Expand Up @@ -753,7 +753,7 @@ static int CeedVectorNorm_Hip(CeedVector vec, CeedNormType type, CeedScalar *nor
// Take reciprocal of a vector on host
//------------------------------------------------------------------------------
static int CeedHostReciprocal_Hip(CeedScalar *h_array, CeedSize length) {
for (CeedSize i = 0; i < length; i++) {
CeedPragmaSIMD for (CeedSize i = 0; i < length; i++) {
if (fabs(h_array[i]) > CEED_EPSILON) h_array[i] = 1. / h_array[i];
}
return CEED_ERROR_SUCCESS;
Expand Down Expand Up @@ -783,7 +783,7 @@ static int CeedVectorReciprocal_Hip(CeedVector vec) {
// Compute x = alpha x on the host
//------------------------------------------------------------------------------
static int CeedHostScale_Hip(CeedScalar *x_array, CeedScalar alpha, CeedSize length) {
for (CeedSize i = 0; i < length; i++) x_array[i] *= alpha;
CeedPragmaSIMD for (CeedSize i = 0; i < length; i++) x_array[i] *= alpha;
return CEED_ERROR_SUCCESS;
}

Expand Down Expand Up @@ -827,11 +827,41 @@ static int CeedVectorScale_Hip(CeedVector x, CeedScalar alpha) {
return CEED_ERROR_SUCCESS;
}

//------------------------------------------------------------------------------
// Filter or clip a vector using a threshold value on the host
//------------------------------------------------------------------------------
static int CeedHostFilter_Hip(CeedScalar *x_array, CeedScalar threshold, CeedSize length) {
CeedPragmaSIMD for (CeedSize i = 0; i < length; i++) {
Comment thread
jeremylt marked this conversation as resolved.
if (fabs(x_array[i]) <= threshold) x_array[i] = 0.0;
}
return CEED_ERROR_SUCCESS;
}

//------------------------------------------------------------------------------
// Filter or clip a vector using a threshold value on device (impl in .cu file)
//------------------------------------------------------------------------------
int CeedDeviceFilter_Hip(CeedScalar *x_array, CeedScalar threshold, CeedSize length);

//------------------------------------------------------------------------------
// Filter or clip a vector using a threshold value
//------------------------------------------------------------------------------
static int CeedVectorFilter_Hip(CeedVector vec, CeedScalar threshold) {
CeedSize length;
CeedVector_Hip *impl;

CeedCallBackend(CeedVectorGetData(vec, &impl));
CeedCallBackend(CeedVectorGetLength(vec, &length));
// Set value for synced device/host array
if (impl->d_array) CeedCallBackend(CeedDeviceFilter_Hip(impl->d_array, threshold, length));
if (impl->h_array) CeedCallBackend(CeedHostFilter_Hip(impl->h_array, threshold, length));
return CEED_ERROR_SUCCESS;
}

//------------------------------------------------------------------------------
// Compute y = alpha x + y on the host
//------------------------------------------------------------------------------
static int CeedHostAXPY_Hip(CeedScalar *y_array, CeedScalar alpha, CeedScalar *x_array, CeedSize length) {
for (CeedSize i = 0; i < length; i++) y_array[i] += alpha * x_array[i];
CeedPragmaSIMD for (CeedSize i = 0; i < length; i++) y_array[i] += alpha * x_array[i];
return CEED_ERROR_SUCCESS;
}

Expand Down Expand Up @@ -881,7 +911,7 @@ static int CeedVectorAXPY_Hip(CeedVector y, CeedScalar alpha, CeedVector x) {
// Compute y = alpha x + beta y on the host
//------------------------------------------------------------------------------
static int CeedHostAXPBY_Hip(CeedScalar *y_array, CeedScalar alpha, CeedScalar beta, CeedScalar *x_array, CeedSize length) {
for (CeedSize i = 0; i < length; i++) y_array[i] = alpha * x_array[i] + beta * y_array[i];
CeedPragmaSIMD for (CeedSize i = 0; i < length; i++) y_array[i] = alpha * x_array[i] + beta * y_array[i];
return CEED_ERROR_SUCCESS;
}

Expand Down Expand Up @@ -916,7 +946,7 @@ static int CeedVectorAXPBY_Hip(CeedVector y, CeedScalar alpha, CeedScalar beta,
// Compute the pointwise multiplication w = x .* y on the host
//------------------------------------------------------------------------------
static int CeedHostPointwiseMult_Hip(CeedScalar *w_array, CeedScalar *x_array, CeedScalar *y_array, CeedSize length) {
for (CeedSize i = 0; i < length; i++) w_array[i] = x_array[i] * y_array[i];
CeedPragmaSIMD for (CeedSize i = 0; i < length; i++) w_array[i] = x_array[i] * y_array[i];
return CEED_ERROR_SUCCESS;
}

Expand Down Expand Up @@ -990,6 +1020,7 @@ int CeedVectorCreate_Hip(CeedSize n, CeedVector vec) {
CeedCallBackend(CeedSetBackendFunction(ceed, "Vector", vec, "Norm", CeedVectorNorm_Hip));
CeedCallBackend(CeedSetBackendFunction(ceed, "Vector", vec, "Reciprocal", CeedVectorReciprocal_Hip));
CeedCallBackend(CeedSetBackendFunction(ceed, "Vector", vec, "Scale", CeedVectorScale_Hip));
CeedCallBackend(CeedSetBackendFunction(ceed, "Vector", vec, "Filter", CeedVectorFilter_Hip));
CeedCallBackend(CeedSetBackendFunction(ceed, "Vector", vec, "AXPY", CeedVectorAXPY_Hip));
CeedCallBackend(CeedSetBackendFunction(ceed, "Vector", vec, "AXPBY", CeedVectorAXPBY_Hip));
CeedCallBackend(CeedSetBackendFunction(ceed, "Vector", vec, "PointwiseMult", CeedVectorPointwiseMult_Hip));
Expand Down
24 changes: 24 additions & 0 deletions backends/hip-ref/kernels/hip-ref-vector.hip.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -124,6 +124,30 @@ extern "C" int CeedDeviceScale_Hip(CeedScalar *x_array, CeedScalar alpha, CeedSi
return 0;
}

//------------------------------------------------------------------------------
// Kernel for filter
//------------------------------------------------------------------------------
__global__ static void filterValueK(CeedScalar *__restrict__ x, CeedScalar threshold, CeedSize size) {
const CeedSize index = threadIdx.x + (CeedSize)blockDim.x * blockIdx.x;

if (index < size) {
if (fabs(x[index]) <= threshold) x[index] = 0.0;
}
}

//------------------------------------------------------------------------------
// Filter or clip vector components to zero if their absolute value is less than or equal to threshold on device
//------------------------------------------------------------------------------
extern "C" int CeedDeviceFilter_Hip(CeedScalar *x_array, CeedScalar threshold, CeedSize length) {
const int block_size = 512;
const CeedSize vec_size = length;
int grid_size = vec_size / block_size;

if (block_size * grid_size < vec_size) grid_size += 1;
hipLaunchKernelGGL(filterValueK, dim3(grid_size), dim3(block_size), 0, 0, x_array, threshold, length);
return 0;
}

//------------------------------------------------------------------------------
// Kernel for axpy
//------------------------------------------------------------------------------
Expand Down
Loading