Skip to content

Updated FIR to add __shared__ optimization of Base_HIP variant. #543

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

Open
mxxw wants to merge 1 commit into
base: develop
Choose a base branch
from
Open

Updated FIR to add __shared__ optimization of Base_HIP variant. #543

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
31 changes: 28 additions & 3 deletions src/apps/FIR-Hip.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -22,8 +22,10 @@ namespace rajaperf
namespace apps
{

#define USE_HIP_CONSTANT_MEMORY
// #undef USE_HIP_CONSTANT_MEMORY
// #define USE_HIP_CONSTANT_MEMORY
#undef USE_HIP_CONSTANT_MEMORY
#define USE_HIP_SHARED_MEMORY
// #undef USE_HIP_SHARED_MEMORY

#if defined(USE_HIP_CONSTANT_MEMORY)

Expand All @@ -47,6 +49,29 @@ __global__ void fir(Real_ptr out, Real_ptr in,
}
}

#elif defined(USE_HIP_SHARED_MEMORY)

__constant__ Real_type coeff[FIR_COEFFLEN];

#define FIR_DATA_SETUP_HIP \
hipErrchk( hipMemcpyToSymbolAsync(HIP_SYMBOL(coeff), coeff_array, FIR_COEFFLEN * sizeof(Real_type), 0, hipMemcpyHostToDevice, res.get_stream()) );


#define FIR_DATA_TEARDOWN_HIP

template < size_t block_size >
__launch_bounds__(block_size)
__global__ void fir(Real_ptr out, Real_ptr in,
const Index_type coefflen,
Index_type iend)
{
__shared__ Real_type coeff_shd[FIR_COEFFLEN];
Index_type i = blockIdx.x * block_size + threadIdx.x;
if (i < iend) {
FIR_BODY_SHD;
}
}

#else // use global memry for coefficients

#define FIR_DATA_SETUP_HIP \
Expand Down Expand Up @@ -99,7 +124,7 @@ void FIR::runHipVariantImpl(VariantID vid)
const size_t grid_size = RAJA_DIVIDE_CEILING_INT(iend, block_size);
constexpr size_t shmem = 0;

#if defined(USE_HIP_CONSTANT_MEMORY)
#if defined(USE_HIP_CONSTANT_MEMORY) || defined(USE_HIP_SHARED_MEMORY)
RPlaunchHipKernel( (fir<block_size>),
grid_size, block_size,
shmem, res.get_stream(),
Expand Down
24 changes: 24 additions & 0 deletions src/apps/FIR.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -51,6 +51,30 @@
} \
out[i] = sum;

#define FIR_BODY_SHD \
Index_type nthds_last_blk = iend & (block_size - 1); /* block_size=2^int */ \
/* case where last block of grid has < FIR_COEFFLEN but > 0 threads */ \
if ( blockIdx.x == gridDim.x - 1 && \
nthds_last_blk < FIR_COEFFLEN && \
nthds_last_blk > 0 ) { \
Index_type ntrips_per_thd = (FIR_COEFFLEN - 1) / nthds_last_blk + 1; \
for (Index_type trip = 0; trip < ntrips_per_thd; trip++) { \
Index_type j = threadIdx.x + trip * nthds_last_blk; \
if ( j < FIR_COEFFLEN ) { \
coeff_shd[j] = coeff[j]; \
} \
} \
} \
else if (threadIdx.x < FIR_COEFFLEN) { \
coeff_shd[threadIdx.x] = coeff[threadIdx.x]; \
} \
__syncthreads(); \
\
Real_type sum = 0.0; \
for (Index_type j = 0; j < coefflen; ++j ) { \
sum += coeff_shd[j]*in[i+j]; \
} \
out[i] = sum;

#include "common/KernelBase.hpp"

Expand Down