feat(simulator): add CUDA and HIP BaseSimulators

GPUSimulators/simulator/__init__.py

@@ -1,4 +1,10 @@
-from .boundary import BoundaryCondition
+from os import environ
 
-# TODO make this dependent on HIP or CUDA
-from .simulator import BaseSimulator
+__env_name = 'GPU_LANG'
+
+if __env_name in environ and environ.get(__env_name).lower() == "cuda":
+    from .cuda_simulator import CudaSimulator as BaseSimulator
+else:
+    from .hip_simulator import HIPSimulator as BaseSimulator
+
+from .boundary import BoundaryCondition

GPUSimulators/simulator/boundary.py

@@ -21,9 +21,10 @@ along with this program.  If not, see <http://www.gnu.org/licenses/>.
 """
 
 # Import packages we need
-import numpy as np
 from enum import IntEnum
 
+import numpy as np
+
 
 class BoundaryCondition(object):
     """

GPUSimulators/simulator/cuda_simulator.py

@@ -0,0 +1,24 @@
+import pycuda.driver as cuda
+
+from GPUSimulators.gpu import KernelContext
+from . import BaseSimulator, BoundaryCondition
+
+
+class CudaSimulator(BaseSimulator):
+    def __init__(self,
+                 context: KernelContext,
+                 nx: int, ny: int,
+                 dx: int, dy: int,
+                 boundary_conditions: BoundaryCondition,
+                 cfl_scale: float,
+                 num_substeps: int,
+                 block_width: int, block_height: int):
+        super().__init__(context, nx, ny, dx, dy, boundary_conditions, cfl_scale, num_substeps, block_width,
+                         block_height)
+
+        # Create a CUDA stream
+        self.stream = cuda.Stream()
+        self.internal_stream = cuda.Stream()
+
+    def synchronize(self):
+        self.stream.synchronize()

GPUSimulators/simulator/hip_simulator.py

@@ -0,0 +1,30 @@
+from hip import hip
+
+from GPUSimulators.common import hip_check
+from GPUSimulators.gpu import KernelContext
+from . import BaseSimulator, BoundaryCondition
+
+
+class HIPSimulator(BaseSimulator):
+    def __init__(self,
+                 context: KernelContext,
+                 nx: int, ny: int,
+                 dx: int, dy: int,
+                 boundary_conditions: BoundaryCondition,
+                 cfl_scale: float,
+                 num_substeps: int,
+                 block_width: int, block_height: int):
+        super().__init__(context, nx, ny, dx, dy, boundary_conditions, cfl_scale, num_substeps, block_width,
+                         block_height)
+
+        # Create a HIP stream
+        self.stream = hip_check(hip.hipStreamCreate())
+        self.internal_stream = hip_check(hip.hipStreamCreate())
+
+    def __del__(self):
+        # Destroy the streams
+        hip_check(hip.hipStreamDestroy(self.stream))
+        hip_check(hip.hipStreamDestroy(self.internal_stream))
+
+    def synchronize(self):
+        hip_check(hip.hipStreamSynchronize(self.stream))

GPUSimulators/simulator/simulator.py

@@ -14,6 +14,7 @@ def get_types(bc):
              'west': BoundaryCondition.Type((bc >> 0) & 0x0000000F)}
     return types
 
+
 class BaseSimulator(object):
 
     def __init__(self,
@@ -71,9 +72,9 @@ class BaseSimulator(object):
             int(np.ceil(self.ny / float(self.block_size[1])))
         )
 
-        # Create a CUDA stream
-        self.stream = cuda.Stream()
-        self.internal_stream = cuda.Stream()
+        # Streams to be implemented in respective language classes
+        self.stream = None
+        self.internal_stream = None
 
         # Keep track of simulation time and number of timesteps
         self.t = 0.0
@@ -143,7 +144,7 @@ class BaseSimulator(object):
         return self.get_output().download(self.stream, variables)
 
     def synchronize(self):
-        self.stream.synchronize()
+        raise NotImplementedError("Needs to be implemented in HIP/CUDA subclass")
 
     def sim_time(self):
         return self.t
@@ -159,21 +160,21 @@ class BaseSimulator(object):
         self.boundary_conditions = boundary_conditions.as_coded_int()
 
     def get_boundary_conditions(self):
-        return BoundaryCondition(get_types())
+        return BoundaryCondition(get_types(self.boundary_conditions))
 
     def substep(self, dt, step_number):
         """
         Function which performs one single substep with stepsize dt
         """
 
-        raise (NotImplementedError("Needs to be implemented in subclass"))
+        raise NotImplementedError("Needs to be implemented in subclass")
 
     def get_output(self):
-        raise (NotImplementedError("Needs to be implemented in subclass"))
+        raise NotImplementedError("Needs to be implemented in subclass")
 
     def check(self):
         self.logger.warning("check() is not implemented - please implement")
         # raise(NotImplementedError("Needs to be implemented in subclass"))
 
     def compute_dt(self):
-        raise (NotImplementedError("Needs to be implemented in subclass"))
+        raise NotImplementedError("Needs to be implemented in subclass")