From 7e05c6f8fd2255751f3e7c0c84d50d068f4e24b3 Mon Sep 17 00:00:00 2001
From: Richard Berger <richard.berger@temple.edu>
Date: Wed, 2 Jun 2021 14:40:37 -0400
Subject: [PATCH 1/3] Update ROCm container defintiion and add missing lib path

---
 tools/singularity/ubuntu20.04_amd_rocm.def | 1 +
 tools/singularity/ubuntu20.04_gpu.def      | 4 ++--
 2 files changed, 3 insertions(+), 2 deletions(-)

diff --git a/tools/singularity/ubuntu20.04_amd_rocm.def b/tools/singularity/ubuntu20.04_amd_rocm.def
index 4eee97ffec..9db8265629 100644
--- a/tools/singularity/ubuntu20.04_amd_rocm.def
+++ b/tools/singularity/ubuntu20.04_amd_rocm.def
@@ -3,6 +3,7 @@ From: ubuntu:20.04
 
 %environment
     export PATH=/usr/lib/ccache:${PATH}:/opt/rocm/bin:/opt/rocm/profiler/bin:/opt/rocm/opencl/bin/x86_64
+    export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:/opt/rocm/lib:/opt/rocm-4.2.0/llvm/lib
 %post
     export DEBIAN_FRONTEND=noninteractive
     apt-get update
diff --git a/tools/singularity/ubuntu20.04_gpu.def b/tools/singularity/ubuntu20.04_gpu.def
index 29938b33f4..44f975d2c8 100644
--- a/tools/singularity/ubuntu20.04_gpu.def
+++ b/tools/singularity/ubuntu20.04_gpu.def
@@ -5,7 +5,7 @@ From: ubuntu:20.04
     export PATH=/usr/lib/ccache:/usr/local/cuda-11.0/bin:${PATH}:/opt/rocm/bin:/opt/rocm/profiler/bin:/opt/rocm/opencl/bin/x86_64
     export CUDADIR=/usr/local/cuda-11.0
     export CUDA_PATH=/usr/local/cuda-11.0
-    export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:/usr/local/cuda-11.0/lib64
+    export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:/usr/local/cuda-11.0/lib64:/opt/rocm/lib:/opt/rocm-4.2.0/llvm/lib
     export LIBRARY_PATH=/usr/local/cuda-11.0/lib64/stubs
 %post
     export DEBIAN_FRONTEND=noninteractive
@@ -131,7 +131,7 @@ From: ubuntu:20.04
     ###########################################################################
 
     export PATH=$PATH:/opt/rocm/bin:/opt/rocm/profiler/bin:/opt/rocm/opencl/bin/x86_64
-    git clone -b rocm-3.7.x https://github.com/ROCmSoftwarePlatform/hipCUB.git
+    git clone -b rocm-4.2.x https://github.com/ROCmSoftwarePlatform/hipCUB.git
     mkdir hipCUB/build
     cd hipCUB/build
     CXX=hipcc cmake -D BUILD_TEST=off ..

From 62423b0c75224a62e08f9bd9f263ce9e281e06a1 Mon Sep 17 00:00:00 2001
From: Richard Berger <richard.berger@temple.edu>
Date: Wed, 2 Jun 2021 15:45:06 -0400
Subject: [PATCH 2/3] Disable new GPU neighbor for HIP

---
 lib/gpu/geryon/hip_device.h | 4 ----
 lib/gpu/lal_neighbor.h      | 4 ++++
 lib/gpu/lal_neighbor_gpu.cu | 4 ++++
 3 files changed, 8 insertions(+), 4 deletions(-)

diff --git a/lib/gpu/geryon/hip_device.h b/lib/gpu/geryon/hip_device.h
index d3917ed0cf..11100cbea1 100644
--- a/lib/gpu/geryon/hip_device.h
+++ b/lib/gpu/geryon/hip_device.h
@@ -8,10 +8,6 @@
 #ifndef HIP_DEVICE
 #define HIP_DEVICE
 
-// workaround after GPU package Feb2021 update
-// todo: make new neighbor code work with HIP
-#define LAL_USE_OLD_NEIGHBOR
-
 #include <hip/hip_runtime.h>
 #include <unordered_map>
 #include <string>
diff --git a/lib/gpu/lal_neighbor.h b/lib/gpu/lal_neighbor.h
index fb854a706c..c1e1a87ef4 100644
--- a/lib/gpu/lal_neighbor.h
+++ b/lib/gpu/lal_neighbor.h
@@ -33,6 +33,10 @@
 #endif
 #endif
 
+#if defined(USE_HIP)
+#define LAL_USE_OLD_NEIGHBOR
+#endif
+
 namespace LAMMPS_AL {
 
 class Neighbor {
diff --git a/lib/gpu/lal_neighbor_gpu.cu b/lib/gpu/lal_neighbor_gpu.cu
index 62c93e6cf1..6fd724b494 100644
--- a/lib/gpu/lal_neighbor_gpu.cu
+++ b/lib/gpu/lal_neighbor_gpu.cu
@@ -40,6 +40,10 @@ _texture_2d( pos_tex,int4);
 #endif
 #endif
 
+#ifdef USE_HIP
+#define LAL_USE_OLD_NEIGHBOR
+#endif
+
 __kernel void calc_cell_id(const numtyp4 *restrict x_,
                            unsigned *restrict cell_id,
                            int *restrict particle_id,

From 8ea9179a00175f4d44c8239c166a679f1d895e81 Mon Sep 17 00:00:00 2001
From: Richard Berger <richard.berger@temple.edu>
Date: Wed, 2 Jun 2021 15:46:57 -0400
Subject: [PATCH 3/3] Correct device pointer datatype in HIP UCL_Const

---
 lib/gpu/geryon/hip_texture.h | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/lib/gpu/geryon/hip_texture.h b/lib/gpu/geryon/hip_texture.h
index 9117adc879..8738f6e2ea 100644
--- a/lib/gpu/geryon/hip_texture.h
+++ b/lib/gpu/geryon/hip_texture.h
@@ -128,12 +128,12 @@ class UCL_Const {
 				    _cq));
   }
   /// Get device ptr associated with object
-  inline const void* begin() const { return &_global; }
+  inline const hipDeviceptr_t * begin() const { return &_global; }
   inline void clear() {}
 
  private:
   hipStream_t _cq;
-  void* _global;
+  hipDeviceptr_t _global;
   size_t _global_bytes;
   friend class UCL_Kernel;
 };