Skip to content

Commit

Permalink
Use getFuserDataSpace with comm
Browse files Browse the repository at this point in the history
  • Loading branch information
@MrBurmark
MrBurmark committed Jan 23, 2024
1 parent 455baee commit f1d0120
Show file tree
Hide file tree
Showing 4 changed files with 84 additions and 84 deletions.
42 changes: 21 additions & 21 deletions src/comm/HALO_EXCHANGE_FUSED-Cuda.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -22,33 +22,33 @@ namespace rajaperf
namespace comm
{

#define HALO_EXCHANGE_FUSED_MANUAL_FUSER_SETUP_CUDA \
#define HALO_EXCHANGE_FUSED_MANUAL_FUSER_SETUP_CUDA(vid) \
Real_ptr* pack_buffer_ptrs; \
Int_ptr* pack_list_ptrs; \
Real_ptr* pack_var_ptrs; \
Index_type* pack_len_ptrs; \
allocData(DataSpace::CudaPinned, pack_buffer_ptrs, num_neighbors * num_vars); \
allocData(DataSpace::CudaPinned, pack_list_ptrs, num_neighbors * num_vars); \
allocData(DataSpace::CudaPinned, pack_var_ptrs, num_neighbors * num_vars); \
allocData(DataSpace::CudaPinned, pack_len_ptrs, num_neighbors * num_vars); \
allocData(getFuserDataSpace(vid), pack_buffer_ptrs, num_neighbors * num_vars); \
allocData(getFuserDataSpace(vid), pack_list_ptrs, num_neighbors * num_vars); \
allocData(getFuserDataSpace(vid), pack_var_ptrs, num_neighbors * num_vars); \
allocData(getFuserDataSpace(vid), pack_len_ptrs, num_neighbors * num_vars); \
Real_ptr* unpack_buffer_ptrs; \
Int_ptr* unpack_list_ptrs; \
Real_ptr* unpack_var_ptrs; \
Index_type* unpack_len_ptrs; \
allocData(DataSpace::CudaPinned, unpack_buffer_ptrs, num_neighbors * num_vars); \
allocData(DataSpace::CudaPinned, unpack_list_ptrs, num_neighbors * num_vars); \
allocData(DataSpace::CudaPinned, unpack_var_ptrs, num_neighbors * num_vars); \
allocData(DataSpace::CudaPinned, unpack_len_ptrs, num_neighbors * num_vars);

#define HALO_EXCHANGE_FUSED_MANUAL_FUSER_TEARDOWN_CUDA \
deallocData(DataSpace::CudaPinned, pack_buffer_ptrs); \
deallocData(DataSpace::CudaPinned, pack_list_ptrs); \
deallocData(DataSpace::CudaPinned, pack_var_ptrs); \
deallocData(DataSpace::CudaPinned, pack_len_ptrs); \
deallocData(DataSpace::CudaPinned, unpack_buffer_ptrs); \
deallocData(DataSpace::CudaPinned, unpack_list_ptrs); \
deallocData(DataSpace::CudaPinned, unpack_var_ptrs); \
deallocData(DataSpace::CudaPinned, unpack_len_ptrs);
allocData(getFuserDataSpace(vid), unpack_buffer_ptrs, num_neighbors * num_vars); \
allocData(getFuserDataSpace(vid), unpack_list_ptrs, num_neighbors * num_vars); \
allocData(getFuserDataSpace(vid), unpack_var_ptrs, num_neighbors * num_vars); \
allocData(getFuserDataSpace(vid), unpack_len_ptrs, num_neighbors * num_vars);

#define HALO_EXCHANGE_FUSED_MANUAL_FUSER_TEARDOWN_CUDA(vid) \
deallocData(getFuserDataSpace(vid), pack_buffer_ptrs); \
deallocData(getFuserDataSpace(vid), pack_list_ptrs); \
deallocData(getFuserDataSpace(vid), pack_var_ptrs); \
deallocData(getFuserDataSpace(vid), pack_len_ptrs); \
deallocData(getFuserDataSpace(vid), unpack_buffer_ptrs); \
deallocData(getFuserDataSpace(vid), unpack_list_ptrs); \
deallocData(getFuserDataSpace(vid), unpack_var_ptrs); \
deallocData(getFuserDataSpace(vid), unpack_len_ptrs);

template < size_t block_size >
__launch_bounds__(block_size)
Expand Down Expand Up @@ -100,7 +100,7 @@ void HALO_EXCHANGE_FUSED::runCudaVariantDirect(VariantID vid)

if ( vid == Base_CUDA ) {

HALO_EXCHANGE_FUSED_MANUAL_FUSER_SETUP_CUDA;
HALO_EXCHANGE_FUSED_MANUAL_FUSER_SETUP_CUDA(Base_CUDA);

startTimer();
for (RepIndex_type irep = 0; irep < run_reps; ++irep) {
Expand Down Expand Up @@ -199,7 +199,7 @@ void HALO_EXCHANGE_FUSED::runCudaVariantDirect(VariantID vid)
}
stopTimer();

HALO_EXCHANGE_FUSED_MANUAL_FUSER_TEARDOWN_CUDA;
HALO_EXCHANGE_FUSED_MANUAL_FUSER_TEARDOWN_CUDA(Base_CUDA);

} else {
getCout() << "\n HALO_EXCHANGE_FUSED : Unknown Cuda variant id = " << vid << std::endl;
Expand Down
42 changes: 21 additions & 21 deletions src/comm/HALO_EXCHANGE_FUSED-Hip.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -22,33 +22,33 @@ namespace rajaperf
namespace comm
{

#define HALO_EXCHANGE_FUSED_MANUAL_FUSER_SETUP_HIP \
#define HALO_EXCHANGE_FUSED_MANUAL_FUSER_SETUP_HIP(vid) \
Real_ptr* pack_buffer_ptrs; \
Int_ptr* pack_list_ptrs; \
Real_ptr* pack_var_ptrs; \
Index_type* pack_len_ptrs; \
allocData(DataSpace::HipPinnedCoarse, pack_buffer_ptrs, num_neighbors * num_vars); \
allocData(DataSpace::HipPinnedCoarse, pack_list_ptrs, num_neighbors * num_vars); \
allocData(DataSpace::HipPinnedCoarse, pack_var_ptrs, num_neighbors * num_vars); \
allocData(DataSpace::HipPinnedCoarse, pack_len_ptrs, num_neighbors * num_vars); \
allocData(getFuserDataSpace(vid), pack_buffer_ptrs, num_neighbors * num_vars); \
allocData(getFuserDataSpace(vid), pack_list_ptrs, num_neighbors * num_vars); \
allocData(getFuserDataSpace(vid), pack_var_ptrs, num_neighbors * num_vars); \
allocData(getFuserDataSpace(vid), pack_len_ptrs, num_neighbors * num_vars); \
Real_ptr* unpack_buffer_ptrs; \
Int_ptr* unpack_list_ptrs; \
Real_ptr* unpack_var_ptrs; \
Index_type* unpack_len_ptrs; \
allocData(DataSpace::HipPinnedCoarse, unpack_buffer_ptrs, num_neighbors * num_vars); \
allocData(DataSpace::HipPinnedCoarse, unpack_list_ptrs, num_neighbors * num_vars); \
allocData(DataSpace::HipPinnedCoarse, unpack_var_ptrs, num_neighbors * num_vars); \
allocData(DataSpace::HipPinnedCoarse, unpack_len_ptrs, num_neighbors * num_vars);

#define HALO_EXCHANGE_FUSED_MANUAL_FUSER_TEARDOWN_HIP \
deallocData(DataSpace::HipPinnedCoarse, pack_buffer_ptrs); \
deallocData(DataSpace::HipPinnedCoarse, pack_list_ptrs); \
deallocData(DataSpace::HipPinnedCoarse, pack_var_ptrs); \
deallocData(DataSpace::HipPinnedCoarse, pack_len_ptrs); \
deallocData(DataSpace::HipPinnedCoarse, unpack_buffer_ptrs); \
deallocData(DataSpace::HipPinnedCoarse, unpack_list_ptrs); \
deallocData(DataSpace::HipPinnedCoarse, unpack_var_ptrs); \
deallocData(DataSpace::HipPinnedCoarse, unpack_len_ptrs);
allocData(getFuserDataSpace(vid), unpack_buffer_ptrs, num_neighbors * num_vars); \
allocData(getFuserDataSpace(vid), unpack_list_ptrs, num_neighbors * num_vars); \
allocData(getFuserDataSpace(vid), unpack_var_ptrs, num_neighbors * num_vars); \
allocData(getFuserDataSpace(vid), unpack_len_ptrs, num_neighbors * num_vars);

#define HALO_EXCHANGE_FUSED_MANUAL_FUSER_TEARDOWN_HIP(vid) \
deallocData(getFuserDataSpace(vid), pack_buffer_ptrs); \
deallocData(getFuserDataSpace(vid), pack_list_ptrs); \
deallocData(getFuserDataSpace(vid), pack_var_ptrs); \
deallocData(getFuserDataSpace(vid), pack_len_ptrs); \
deallocData(getFuserDataSpace(vid), unpack_buffer_ptrs); \
deallocData(getFuserDataSpace(vid), unpack_list_ptrs); \
deallocData(getFuserDataSpace(vid), unpack_var_ptrs); \
deallocData(getFuserDataSpace(vid), unpack_len_ptrs);

template < size_t block_size >
__launch_bounds__(block_size)
Expand Down Expand Up @@ -100,7 +100,7 @@ void HALO_EXCHANGE_FUSED::runHipVariantDirect(VariantID vid)

if ( vid == Base_HIP ) {

HALO_EXCHANGE_FUSED_MANUAL_FUSER_SETUP_HIP;
HALO_EXCHANGE_FUSED_MANUAL_FUSER_SETUP_HIP(Base_HIP);

startTimer();
for (RepIndex_type irep = 0; irep < run_reps; ++irep) {
Expand Down Expand Up @@ -199,7 +199,7 @@ void HALO_EXCHANGE_FUSED::runHipVariantDirect(VariantID vid)
}
stopTimer();

HALO_EXCHANGE_FUSED_MANUAL_FUSER_TEARDOWN_HIP;
HALO_EXCHANGE_FUSED_MANUAL_FUSER_TEARDOWN_HIP(Base_HIP);

} else {
getCout() << "\n HALO_EXCHANGE_FUSED : Unknown Hip variant id = " << vid << std::endl;
Expand Down
42 changes: 21 additions & 21 deletions src/comm/HALO_PACKING_FUSED-Cuda.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -22,33 +22,33 @@ namespace rajaperf
namespace comm
{

#define HALO_PACKING_FUSED_MANUAL_FUSER_SETUP_CUDA \
#define HALO_PACKING_FUSED_MANUAL_FUSER_SETUP_CUDA(vid) \
Real_ptr* pack_buffer_ptrs; \
Int_ptr* pack_list_ptrs; \
Real_ptr* pack_var_ptrs; \
Index_type* pack_len_ptrs; \
allocData(DataSpace::CudaPinned, pack_buffer_ptrs, num_neighbors * num_vars); \
allocData(DataSpace::CudaPinned, pack_list_ptrs, num_neighbors * num_vars); \
allocData(DataSpace::CudaPinned, pack_var_ptrs, num_neighbors * num_vars); \
allocData(DataSpace::CudaPinned, pack_len_ptrs, num_neighbors * num_vars); \
allocData(getFuserDataSpace(vid), pack_buffer_ptrs, num_neighbors * num_vars); \
allocData(getFuserDataSpace(vid), pack_list_ptrs, num_neighbors * num_vars); \
allocData(getFuserDataSpace(vid), pack_var_ptrs, num_neighbors * num_vars); \
allocData(getFuserDataSpace(vid), pack_len_ptrs, num_neighbors * num_vars); \
Real_ptr* unpack_buffer_ptrs; \
Int_ptr* unpack_list_ptrs; \
Real_ptr* unpack_var_ptrs; \
Index_type* unpack_len_ptrs; \
allocData(DataSpace::CudaPinned, unpack_buffer_ptrs, num_neighbors * num_vars); \
allocData(DataSpace::CudaPinned, unpack_list_ptrs, num_neighbors * num_vars); \
allocData(DataSpace::CudaPinned, unpack_var_ptrs, num_neighbors * num_vars); \
allocData(DataSpace::CudaPinned, unpack_len_ptrs, num_neighbors * num_vars);

#define HALO_PACKING_FUSED_MANUAL_FUSER_TEARDOWN_CUDA \
deallocData(DataSpace::CudaPinned, pack_buffer_ptrs); \
deallocData(DataSpace::CudaPinned, pack_list_ptrs); \
deallocData(DataSpace::CudaPinned, pack_var_ptrs); \
deallocData(DataSpace::CudaPinned, pack_len_ptrs); \
deallocData(DataSpace::CudaPinned, unpack_buffer_ptrs); \
deallocData(DataSpace::CudaPinned, unpack_list_ptrs); \
deallocData(DataSpace::CudaPinned, unpack_var_ptrs); \
deallocData(DataSpace::CudaPinned, unpack_len_ptrs);
allocData(getFuserDataSpace(vid), unpack_buffer_ptrs, num_neighbors * num_vars); \
allocData(getFuserDataSpace(vid), unpack_list_ptrs, num_neighbors * num_vars); \
allocData(getFuserDataSpace(vid), unpack_var_ptrs, num_neighbors * num_vars); \
allocData(getFuserDataSpace(vid), unpack_len_ptrs, num_neighbors * num_vars);

#define HALO_PACKING_FUSED_MANUAL_FUSER_TEARDOWN_CUDA(vid) \
deallocData(getFuserDataSpace(vid), pack_buffer_ptrs); \
deallocData(getFuserDataSpace(vid), pack_list_ptrs); \
deallocData(getFuserDataSpace(vid), pack_var_ptrs); \
deallocData(getFuserDataSpace(vid), pack_len_ptrs); \
deallocData(getFuserDataSpace(vid), unpack_buffer_ptrs); \
deallocData(getFuserDataSpace(vid), unpack_list_ptrs); \
deallocData(getFuserDataSpace(vid), unpack_var_ptrs); \
deallocData(getFuserDataSpace(vid), unpack_len_ptrs);

template < size_t block_size >
__launch_bounds__(block_size)
Expand Down Expand Up @@ -104,7 +104,7 @@ void HALO_PACKING_FUSED::runCudaVariantDirect(VariantID vid)

if ( vid == Base_CUDA ) {

HALO_PACKING_FUSED_MANUAL_FUSER_SETUP_CUDA;
HALO_PACKING_FUSED_MANUAL_FUSER_SETUP_CUDA(Base_CUDA);

startTimer();
for (RepIndex_type irep = 0; irep < run_reps; ++irep) {
Expand Down Expand Up @@ -190,7 +190,7 @@ void HALO_PACKING_FUSED::runCudaVariantDirect(VariantID vid)
}
stopTimer();

HALO_PACKING_FUSED_MANUAL_FUSER_TEARDOWN_CUDA;
HALO_PACKING_FUSED_MANUAL_FUSER_TEARDOWN_CUDA(Base_CUDA);

} else {
getCout() << "\n HALO_PACKING_FUSED : Unknown Cuda variant id = " << vid << std::endl;
Expand Down
42 changes: 21 additions & 21 deletions src/comm/HALO_PACKING_FUSED-Hip.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -22,33 +22,33 @@ namespace rajaperf
namespace comm
{

#define HALO_PACKING_FUSED_MANUAL_FUSER_SETUP_HIP \
#define HALO_PACKING_FUSED_MANUAL_FUSER_SETUP_HIP(vid) \
Real_ptr* pack_buffer_ptrs; \
Int_ptr* pack_list_ptrs; \
Real_ptr* pack_var_ptrs; \
Index_type* pack_len_ptrs; \
allocData(DataSpace::HipPinnedCoarse, pack_buffer_ptrs, num_neighbors * num_vars); \
allocData(DataSpace::HipPinnedCoarse, pack_list_ptrs, num_neighbors * num_vars); \
allocData(DataSpace::HipPinnedCoarse, pack_var_ptrs, num_neighbors * num_vars); \
allocData(DataSpace::HipPinnedCoarse, pack_len_ptrs, num_neighbors * num_vars); \
allocData(getFuserDataSpace(vid), pack_buffer_ptrs, num_neighbors * num_vars); \
allocData(getFuserDataSpace(vid), pack_list_ptrs, num_neighbors * num_vars); \
allocData(getFuserDataSpace(vid), pack_var_ptrs, num_neighbors * num_vars); \
allocData(getFuserDataSpace(vid), pack_len_ptrs, num_neighbors * num_vars); \
Real_ptr* unpack_buffer_ptrs; \
Int_ptr* unpack_list_ptrs; \
Real_ptr* unpack_var_ptrs; \
Index_type* unpack_len_ptrs; \
allocData(DataSpace::HipPinnedCoarse, unpack_buffer_ptrs, num_neighbors * num_vars); \
allocData(DataSpace::HipPinnedCoarse, unpack_list_ptrs, num_neighbors * num_vars); \
allocData(DataSpace::HipPinnedCoarse, unpack_var_ptrs, num_neighbors * num_vars); \
allocData(DataSpace::HipPinnedCoarse, unpack_len_ptrs, num_neighbors * num_vars);

#define HALO_PACKING_FUSED_MANUAL_FUSER_TEARDOWN_HIP \
deallocData(DataSpace::HipPinnedCoarse, pack_buffer_ptrs); \
deallocData(DataSpace::HipPinnedCoarse, pack_list_ptrs); \
deallocData(DataSpace::HipPinnedCoarse, pack_var_ptrs); \
deallocData(DataSpace::HipPinnedCoarse, pack_len_ptrs); \
deallocData(DataSpace::HipPinnedCoarse, unpack_buffer_ptrs); \
deallocData(DataSpace::HipPinnedCoarse, unpack_list_ptrs); \
deallocData(DataSpace::HipPinnedCoarse, unpack_var_ptrs); \
deallocData(DataSpace::HipPinnedCoarse, unpack_len_ptrs);
allocData(getFuserDataSpace(vid), unpack_buffer_ptrs, num_neighbors * num_vars); \
allocData(getFuserDataSpace(vid), unpack_list_ptrs, num_neighbors * num_vars); \
allocData(getFuserDataSpace(vid), unpack_var_ptrs, num_neighbors * num_vars); \
allocData(getFuserDataSpace(vid), unpack_len_ptrs, num_neighbors * num_vars);

#define HALO_PACKING_FUSED_MANUAL_FUSER_TEARDOWN_HIP(vid) \
deallocData(getFuserDataSpace(vid), pack_buffer_ptrs); \
deallocData(getFuserDataSpace(vid), pack_list_ptrs); \
deallocData(getFuserDataSpace(vid), pack_var_ptrs); \
deallocData(getFuserDataSpace(vid), pack_len_ptrs); \
deallocData(getFuserDataSpace(vid), unpack_buffer_ptrs); \
deallocData(getFuserDataSpace(vid), unpack_list_ptrs); \
deallocData(getFuserDataSpace(vid), unpack_var_ptrs); \
deallocData(getFuserDataSpace(vid), unpack_len_ptrs);

template < size_t block_size >
__launch_bounds__(block_size)
Expand Down Expand Up @@ -104,7 +104,7 @@ void HALO_PACKING_FUSED::runHipVariantDirect(VariantID vid)

if ( vid == Base_HIP ) {

HALO_PACKING_FUSED_MANUAL_FUSER_SETUP_HIP;
HALO_PACKING_FUSED_MANUAL_FUSER_SETUP_HIP(Base_HIP);

startTimer();
for (RepIndex_type irep = 0; irep < run_reps; ++irep) {
Expand Down Expand Up @@ -190,7 +190,7 @@ void HALO_PACKING_FUSED::runHipVariantDirect(VariantID vid)
}
stopTimer();

HALO_PACKING_FUSED_MANUAL_FUSER_TEARDOWN_HIP;
HALO_PACKING_FUSED_MANUAL_FUSER_TEARDOWN_HIP(Base_HIP);

} else {
getCout() << "\n HALO_PACKING_FUSED : Unknown Hip variant id = " << vid << std::endl;
Expand Down

0 comments on commit f1d0120

Please sign in to comment.