From da9c12015d8618a8331c091b6683aea82d9db1ba Mon Sep 17 00:00:00 2001
From: "Kurt A. O'Hearn" <ohearnku@msu.edu>
Date: Wed, 31 May 2023 13:36:48 -0400
Subject: [PATCH] Fixes from backporting CUDA changes to HIP.
---
PG-PuReMD/src/hip/hip_forces.cu | 6 +++---
PG-PuReMD/src/hip/hip_forces.h | 6 +++---
PG-PuReMD/src/hip/hip_system_props.cu | 28 +++++++++++++--------------
3 files changed, 20 insertions(+), 20 deletions(-)
diff --git a/PG-PuReMD/src/hip/hip_forces.cu b/PG-PuReMD/src/hip/hip_forces.cu
index 1818518..e654701 100644
--- a/PG-PuReMD/src/hip/hip_forces.cu
+++ b/PG-PuReMD/src/hip/hip_forces.cu
@@ -1921,7 +1921,7 @@ void Hip_Init_Neighbor_Indices( reax_system * const system,
*
* system: atomic system info. */
void Hip_Init_HBond_Indices( reax_system * const system, storage * const workspace,
- reax_list * const hbond_list, int block_size, cudaStream_t s )
+ reax_list * const hbond_list, int block_size, hipStream_t s )
{
int blocks, *temp;
@@ -1946,7 +1946,7 @@ void Hip_Init_HBond_Indices( reax_system * const system, storage * const workspa
*
* system: atomic system info. */
void Hip_Init_Bond_Indices( reax_system * const system, reax_list * const bond_list,
- int block_size, cudaStream_t s )
+ int block_size, hipStream_t s )
{
int blocks;
@@ -1969,7 +1969,7 @@ void Hip_Init_Bond_Indices( reax_system * const system, reax_list * const bond_l
* system: atomic system info.
* H: charge matrix */
void Hip_Init_Sparse_Matrix_Indices( reax_system * const system, sparse_matrix * const H,
- int block_size, cudaStream_t s )
+ int block_size, hipStream_t s )
{
int blocks;
diff --git a/PG-PuReMD/src/hip/hip_forces.h b/PG-PuReMD/src/hip/hip_forces.h
index 3acd766..81a1d87 100644
--- a/PG-PuReMD/src/hip/hip_forces.h
+++ b/PG-PuReMD/src/hip/hip_forces.h
@@ -9,12 +9,12 @@ void Hip_Init_Neighbor_Indices( reax_system * const, control_params const * cons
reax_list * const );
void Hip_Init_HBond_Indices( reax_system * const, storage * const, reax_list * const,
- int, cudaStream_t );
+ int, hipStream_t );
-void Hip_Init_Bond_Indices( reax_system * const, reax_list * const, int, cudaStream_t );
+void Hip_Init_Bond_Indices( reax_system * const, reax_list * const, int, hipStream_t );
void Hip_Init_Sparse_Matrix_Indices( reax_system * const, sparse_matrix * const, int,
- cudaStream_t );
+ hipStream_t );
void Hip_Init_Three_Body_Indices( int *, int, reax_list ** );
diff --git a/PG-PuReMD/src/hip/hip_system_props.cu b/PG-PuReMD/src/hip/hip_system_props.cu
index f0c94aa..9433a13 100644
--- a/PG-PuReMD/src/hip/hip_system_props.cu
+++ b/PG-PuReMD/src/hip/hip_system_props.cu
@@ -148,7 +148,7 @@ GPU_GLOBAL void k_center_of_mass_amcm( single_body_parameters *sbp,
GPU_GLOBAL void k_compute_inertial_tensor_part1( single_body_parameters *sbp,
reax_atom *atoms, real *t_g, real xcm0, real xcm1, real xcm2, size_t n )
{
- extern __shared__ hipcub::BlockReduce<double, DEF_BLOCK_SIZE>::TempStorage temp_block[];
+ extern __shared__ hipcub::BlockReduce<double, GPU_BLOCK_SIZE>::TempStorage temp_block[];
unsigned int i;
real xx, xy, xz, yy, yz, zz, m;
rvec diff, xcm;
@@ -179,17 +179,17 @@ GPU_GLOBAL void k_compute_inertial_tensor_part1( single_body_parameters *sbp,
zz = 0.0;
}
- xx = hipcub::BlockReduce<double, DEF_BLOCK_SIZE>(*temp_block).Sum(xx);
+ xx = hipcub::BlockReduce<double, GPU_BLOCK_SIZE>(*temp_block).Sum(xx);
__syncthreads( );
- xy = hipcub::BlockReduce<double, DEF_BLOCK_SIZE>(*temp_block).Sum(xy);
+ xy = hipcub::BlockReduce<double, GPU_BLOCK_SIZE>(*temp_block).Sum(xy);
__syncthreads( );
- xz = hipcub::BlockReduce<double, DEF_BLOCK_SIZE>(*temp_block).Sum(xz);
+ xz = hipcub::BlockReduce<double, GPU_BLOCK_SIZE>(*temp_block).Sum(xz);
__syncthreads( );
- yy = hipcub::BlockReduce<double, DEF_BLOCK_SIZE>(*temp_block).Sum(yy);
+ yy = hipcub::BlockReduce<double, GPU_BLOCK_SIZE>(*temp_block).Sum(yy);
__syncthreads( );
- yz = hipcub::BlockReduce<double, DEF_BLOCK_SIZE>(*temp_block).Sum(yz);
+ yz = hipcub::BlockReduce<double, GPU_BLOCK_SIZE>(*temp_block).Sum(yz);
__syncthreads( );
- zz = hipcub::BlockReduce<double, DEF_BLOCK_SIZE>(*temp_block).Sum(zz);
+ zz = hipcub::BlockReduce<double, GPU_BLOCK_SIZE>(*temp_block).Sum(zz);
/* one thread writes the block-level partial sum
* of the reduction back to global memory */
@@ -207,7 +207,7 @@ GPU_GLOBAL void k_compute_inertial_tensor_part1( single_body_parameters *sbp,
GPU_GLOBAL void k_compute_inertial_tensor_part2( real *input, real *output, size_t n )
{
- extern __shared__ hipcub::BlockReduce<double, DEF_BLOCK_SIZE>::TempStorage temp_block[];
+ extern __shared__ hipcub::BlockReduce<double, GPU_BLOCK_SIZE>::TempStorage temp_block[];
unsigned int i;
real xx, xy, xz, yy, yz, zz;
@@ -232,17 +232,17 @@ GPU_GLOBAL void k_compute_inertial_tensor_part2( real *input, real *output, size
zz = 0.0;
}
- xx = hipcub::BlockReduce<double, DEF_BLOCK_SIZE>(*temp_block).Sum(xx);
+ xx = hipcub::BlockReduce<double, GPU_BLOCK_SIZE>(*temp_block).Sum(xx);
__syncthreads( );
- xy = hipcub::BlockReduce<double, DEF_BLOCK_SIZE>(*temp_block).Sum(xy);
+ xy = hipcub::BlockReduce<double, GPU_BLOCK_SIZE>(*temp_block).Sum(xy);
__syncthreads( );
- xz = hipcub::BlockReduce<double, DEF_BLOCK_SIZE>(*temp_block).Sum(xz);
+ xz = hipcub::BlockReduce<double, GPU_BLOCK_SIZE>(*temp_block).Sum(xz);
__syncthreads( );
- yy = hipcub::BlockReduce<double, DEF_BLOCK_SIZE>(*temp_block).Sum(yy);
+ yy = hipcub::BlockReduce<double, GPU_BLOCK_SIZE>(*temp_block).Sum(yy);
__syncthreads( );
- yz = hipcub::BlockReduce<double, DEF_BLOCK_SIZE>(*temp_block).Sum(yz);
+ yz = hipcub::BlockReduce<double, GPU_BLOCK_SIZE>(*temp_block).Sum(yz);
__syncthreads( );
- zz = hipcub::BlockReduce<double, DEF_BLOCK_SIZE>(*temp_block).Sum(zz);
+ zz = hipcub::BlockReduce<double, GPU_BLOCK_SIZE>(*temp_block).Sum(zz);
if ( threadIdx.x == 0 )
{
--
GitLab