From 0325047c01dac6bdafc14ea16e1e34d439ec55a7 Mon Sep 17 00:00:00 2001
From: Axel Kohlmeyer <akohlmey@gmail.com>
Date: Thu, 21 Oct 2021 07:33:00 -0400
Subject: [PATCH] update a few GPU kernels so they can be compiled on GPUs
 without double precisions support

---
 lib/gpu/lal_atom.cu |  2 +-
 lib/gpu/lal_atom.h  |  2 +-
 lib/gpu/lal_zbl.cu  | 12 ++++++------
 lib/gpu/lal_zbl.h   |  6 +++---
 4 files changed, 11 insertions(+), 11 deletions(-)

diff --git a/lib/gpu/lal_atom.cu b/lib/gpu/lal_atom.cu
index 1418459301..287d72803c 100644
--- a/lib/gpu/lal_atom.cu
+++ b/lib/gpu/lal_atom.cu
@@ -18,7 +18,7 @@
 #endif
 
 __kernel void kernel_cast_x(__global numtyp4 *restrict x_type,
-                            const __global double *restrict x,
+                            const __global numtyp *restrict x,
                             const __global int *restrict type,
                             const int nall) {
   int ii=GLOBAL_ID_X;
diff --git a/lib/gpu/lal_atom.h b/lib/gpu/lal_atom.h
index 3cf97d94a0..77c1faa784 100644
--- a/lib/gpu/lal_atom.h
+++ b/lib/gpu/lal_atom.h
@@ -475,7 +475,7 @@ class Atom {
   UCL_Vector<numtyp,numtyp> v;
 
   #ifdef GPU_CAST
-  UCL_Vector<double,double> x_cast;
+  UCL_Vector<numtyp,numtyp> x_cast;
   UCL_Vector<int,int> type_cast;
   #endif
 
diff --git a/lib/gpu/lal_zbl.cu b/lib/gpu/lal_zbl.cu
index 09e1b4f6bb..2a7d4795da 100644
--- a/lib/gpu/lal_zbl.cu
+++ b/lib/gpu/lal_zbl.cu
@@ -82,9 +82,9 @@ __kernel void k_zbl(const __global numtyp4 *restrict x_,
                     const __global numtyp4 *restrict coeff1,
                     const __global numtyp4 *restrict coeff2,
                     const __global numtyp4 *restrict coeff3,
-                    const double cut_globalsq,
-                    const double cut_innersq,
-                    const double cut_inner,
+                    const numtyp cut_globalsq,
+                    const numtyp cut_innersq,
+                    const numtyp cut_inner,
                     const int lj_types,
                     const __global int *dev_nbor,
                     const __global int *dev_packed,
@@ -174,9 +174,9 @@ __kernel void k_zbl_fast(const __global numtyp4 *restrict x_,
                          const __global numtyp4 *restrict coeff1_in,
                          const __global numtyp4 *restrict coeff2_in,
                          const __global numtyp4 *restrict coeff3_in,
-                         const double cut_globalsq,
-                         const double cut_innersq,
-                         const double cut_inner,
+                         const numtyp cut_globalsq,
+                         const numtyp cut_innersq,
+                         const numtyp cut_inner,
                          const __global int *dev_nbor,
                          const __global int *dev_packed,
                          __global acctyp4 *restrict ans,
diff --git a/lib/gpu/lal_zbl.h b/lib/gpu/lal_zbl.h
index af4f1b2eac..b7b525661c 100644
--- a/lib/gpu/lal_zbl.h
+++ b/lib/gpu/lal_zbl.h
@@ -67,9 +67,9 @@ class ZBL : public BaseAtomic<numtyp, acctyp> {
   /// If atom type constants fit in shared memory, use fast kernels
   bool shared_types;
 
-  double _cut_globalsq;
-  double _cut_innersq;
-  double _cut_inner;
+  numtyp _cut_globalsq;
+  numtyp _cut_innersq;
+  numtyp _cut_inner;
 
   /// Number of atom types
   int _lj_types;