diff --git a/backends/cuda-ref/ceed-cuda-ref-vector.c b/backends/cuda-ref/ceed-cuda-ref-vector.c index de66e1cc16..6feb0f98b4 100644 --- a/backends/cuda-ref/ceed-cuda-ref-vector.c +++ b/backends/cuda-ref/ceed-cuda-ref-vector.c @@ -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; } @@ -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; } @@ -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; } @@ -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; @@ -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; } @@ -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; } @@ -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; } @@ -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; } diff --git a/backends/hip-ref/ceed-hip-ref-vector.c b/backends/hip-ref/ceed-hip-ref-vector.c index 667555db69..e444b0c5a8 100644 --- a/backends/hip-ref/ceed-hip-ref-vector.c +++ b/backends/hip-ref/ceed-hip-ref-vector.c @@ -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) { - 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; } @@ -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; } @@ -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; @@ -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; } @@ -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++) { + 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; } @@ -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; } @@ -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; } @@ -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)); diff --git a/backends/hip-ref/kernels/hip-ref-vector.hip.cpp b/backends/hip-ref/kernels/hip-ref-vector.hip.cpp index b9f81032b5..95df8a0913 100644 --- a/backends/hip-ref/kernels/hip-ref-vector.hip.cpp +++ b/backends/hip-ref/kernels/hip-ref-vector.hip.cpp @@ -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 //------------------------------------------------------------------------------