potential minimum completed

kernels/potential_evaluated_in_phi02_phi01_plane.py

@@ -0,0 +1,76 @@
+"""Potential is plotted for
+- Fixed Phi L
+- Fixed Phi R
+- Minimised phi03
+
+and for a range of phi02-phi01 values
+"""
+
+from typing import List, Callable, Tuple, Optional
+
+from numba import cuda
+from numba.cuda.cudadrv.devicearray import DeviceNDArray
+import numpy as np
+
+
+class PotentialEvaluatedInPhi02Phi02Plane:
+    def __init__(
+        self, number_of_phi_points: int, potential_function_cuda: Callable,
+    ):
+        self.NUMBER_OF_PHI_POINTS = number_of_phi_points
+        self.potential_function_cuda = potential_function_cuda
+
+        self.kernel = self.kernel_wrapper()
+
+    def kernel_wrapper(self):
+        NUMBER_OF_PHI_POINTS = self.NUMBER_OF_PHI_POINTS
+        potential_function_cuda = self.potential_function_cuda
+
+        @cuda.jit
+        def kernel(
+            phi01_array: List[float],
+            phi02_array: List[float],
+            phi03_array: List[float],
+            phi_l: float,
+            phi_r: float,
+            alpha: float,
+            array_out: DeviceNDArray,
+        ):
+            """
+            phixx_array:        array of the values for phi03 - float32
+            phi_l, phi_r:       fixed values for the externally applied phase
+            alpha:              variables parameter
+            array_out:          allocate either with cuda.device_array or passing in a numpy array - floa32
+            """
+
+            phi01_idx = cuda.blockIdx.x
+            phi02_idx = cuda.blockIdx.y
+            phi03_idx = cuda.threadIdx.x
+            potential_at_phi03 = cuda.shared.array(
+                shape=(NUMBER_OF_PHI_POINTS), dtype=np.float32
+            )
+
+            # Traverse over all the phi03 values
+            while phi03_idx < NUMBER_OF_PHI_POINTS:
+                potential_at_phi03[phi03_idx] = potential_function_cuda(
+                    (
+                        phi01_array[phi01_idx],
+                        phi02_array[phi02_idx],
+                        phi03_array[phi03_idx],
+                    ),
+                    phi_l,
+                    phi_r,
+                    alpha,
+                )
+                phi03_idx += cuda.blockDim.x
+            cuda.syncthreads()
+
+            # Then step through and find minimal value
+            min_potential = potential_at_phi03[0]
+            for potential in potential_at_phi03[1:]:
+                if potential < min_potential:
+                    min_potential = potential
+
+            array_out[phi01_idx][phi02_idx] = min_potential
+
+        return kernel

kernels/potential_evaluator.py

@@ -14,55 +14,6 @@ def __init__(
         self.potential_function_cuda = potential_function_cuda
 
         self.kernel = self.kernel_wrapper()
-        self.gpu_info = gpu_check()
-
-    def allocate_max_threads(
-        self, user_defined_number: Optional[int] = None, verbose=False
-    ) -> Tuple[int, int, int]:
-        if verbose:
-            print(
-                f"""Thread parameters:
-        > Max threads per block: {self.gpu_info['max_threads_per_block']}
-        > Max threads in x: {self.gpu_info['max_block_dim_x']}
-        > Max threads in y: {self.gpu_info['max_block_dim_y']}
-        > Max threads in z: {self.gpu_info['max_block_dim_z']}"""
-            )
-        max_threads_approximation = int(
-            self.gpu_info["max_threads_per_block"] ** (1 / 3)
-        )
-        if user_defined_number is not None:
-            max_threads_approximation = user_defined_number
-
-        max_thread_allocation = (
-            min(max_threads_approximation, self.gpu_info["max_block_dim_x"]),
-            min(max_threads_approximation, self.gpu_info["max_block_dim_y"]),
-            min(max_threads_approximation, self.gpu_info["max_block_dim_z"]),
-        )
-        print(f"🐳 {'Allocating':<20} THREADS_PER_BLOCK = {max_thread_allocation}")
-
-        return max_thread_allocation
-
-    def verify_blocks_per_grid(self, blocks_per_grid: Tuple, verbose=False) -> bool:
-        if verbose:
-            print(
-                f"""Block parameters:
-        > Max blocks in x: {self.gpu_info['max_grid_dim_x']}
-        > Max blocks in y: {self.gpu_info['max_grid_dim_y']}
-        > Max blocks in z: {self.gpu_info['max_grid_dim_z']}"""
-            )
-        for (block_dim, max_dim) in zip(
-            blocks_per_grid,
-            [
-                self.gpu_info["max_grid_dim_x"],
-                self.gpu_info["max_grid_dim_y"],
-                self.gpu_info["max_grid_dim_z"],
-            ],
-        ):
-            if block_dim > max_dim:
-                print("🦑 Allocating too many blocks")
-                return False
-        print(f"🐳 {'Verified':<20} BLOCKS_PER_GRID={blocks_per_grid}")
-        return True
 
     def kernel_wrapper(self):
         NUMBER_OF_PHI_POINTS = self.NUMBER_OF_PHI_POINTS

tests/test_potential_evaluated_in_phi02_phi01_plane.py

@@ -0,0 +1,51 @@
+import unittest
+from unittest.mock import Mock
+from unittest.mock import patch
+
+from kernels.potential_evaluated_in_phi02_phi01_plane import (
+    PotentialEvaluatedInPhi02Phi02Plane,
+)
+from template import mock_potential_function_cuda
+import numpy as np
+from numba import cuda
+
+
+class TestPotentialEvaluatedInPhi02Phi02Plane(unittest.TestCase):
+    def setUp(self):
+
+        self.NUMBER_OF_PHI_POINTS = 3
+
+        self.sut = PotentialEvaluatedInPhi02Phi02Plane(
+            self.NUMBER_OF_PHI_POINTS, mock_potential_function_cuda,
+        )
+
+    def tearDown(self):
+        pass
+
+    def test(self):
+        phi01_array = np.array([4, 5, 6])
+        phi02_array = np.array([4, 5, 6])
+        phi03_array = np.array([4, 5, 6])
+        alpha = 7
+        DEVICE_out_array = cuda.device_array(
+            shape=(self.NUMBER_OF_PHI_POINTS, self.NUMBER_OF_PHI_POINTS,),
+            dtype=np.float32,
+        )
+
+        self.sut.kernel[(self.NUMBER_OF_PHI_POINTS, self.NUMBER_OF_PHI_POINTS), 10,](
+            cuda.to_device(phi01_array),
+            cuda.to_device(phi01_array),
+            cuda.to_device(phi01_array),
+            0,
+            0,
+            alpha,
+            DEVICE_out_array,
+        )
+        expected = np.array(
+            [
+                [744400, 745400, 746400],
+                [754400, 755400, 756400],
+                [764400, 765400, 766400],
+            ]
+        )
+        np.all(expected == DEVICE_out_array.copy_to_host())

tests/test_potential_evaluator.py

@@ -4,6 +4,7 @@
 
 from kernels.potential_evaluator import PotentialEvaluator
 from template import mock_potential_function_cuda
+from utils.info import allocate_max_threads
 import numpy as np
 from numba import cuda
 
@@ -37,7 +38,7 @@ def tearDown(self):
         pass
 
     def test(self):
-        THREADS_PER_BLOCK = self.sut.allocate_max_threads()
+        THREADS_PER_BLOCK = allocate_max_threads()
 
         self.sut.kernel[
             (self.NUMBER_OF_FIELD_POINTS_PER_RUN, self.NUMBER_OF_FIELD_POINTS_PER_RUN),
@@ -76,7 +77,7 @@ def test(self):
         np.all(expected == self.DEVICE_out_array.copy_to_host())
 
     def test_offset(self):
-        THREADS_PER_BLOCK = self.sut.allocate_max_threads()
+        THREADS_PER_BLOCK = allocate_max_threads()
 
         self.sut.kernel[
             (self.NUMBER_OF_FIELD_POINTS_PER_RUN, self.NUMBER_OF_FIELD_POINTS_PER_RUN),

utils/info.py

@@ -1,3 +1,5 @@
+from typing import List, Callable, Tuple, Optional
+
 import numba as nb
 from numba import cuda
 
@@ -26,3 +28,54 @@ def gpu_check():
         # )
         return parameters
     raise RuntimeError("Missing GPU")
+
+
+def allocate_max_threads(
+    user_defined_number: Optional[int] = None, verbose=False
+) -> Tuple[int, int, int]:
+    gpu_info = gpu_check()
+    if verbose:
+        print(
+            f"""Thread parameters:
+    > Max threads per block: {gpu_info['max_threads_per_block']}
+    > Max threads in x: {gpu_info['max_block_dim_x']}
+    > Max threads in y: {gpu_info['max_block_dim_y']}
+    > Max threads in z: {gpu_info['max_block_dim_z']}"""
+        )
+    max_threads_approximation = int(gpu_info["max_threads_per_block"] ** (1 / 3))
+    if user_defined_number is not None:
+        max_threads_approximation = user_defined_number
+
+    max_thread_allocation = (
+        min(max_threads_approximation, gpu_info["max_block_dim_x"]),
+        min(max_threads_approximation, gpu_info["max_block_dim_y"]),
+        min(max_threads_approximation, gpu_info["max_block_dim_z"]),
+    )
+    print(f"🐳 {'Allocating':<20} THREADS_PER_BLOCK = {max_thread_allocation}")
+
+    return max_thread_allocation
+
+
+def verify_blocks_per_grid(blocks_per_grid: Tuple, verbose=False) -> bool:
+    gpu_info = gpu_check()
+
+    if verbose:
+        print(
+            f"""Block parameters:
+    > Max blocks in x: {gpu_info['max_grid_dim_x']}
+    > Max blocks in y: {gpu_info['max_grid_dim_y']}
+    > Max blocks in z: {gpu_info['max_grid_dim_z']}"""
+        )
+    for (block_dim, max_dim) in zip(
+        blocks_per_grid,
+        [
+            gpu_info["max_grid_dim_x"],
+            gpu_info["max_grid_dim_y"],
+            gpu_info["max_grid_dim_z"],
+        ],
+    ):
+        if block_dim > max_dim:
+            print("🦑 Allocating too many blocks")
+            return False
+    print(f"🐳 {'Verified':<20} BLOCKS_PER_GRID={blocks_per_grid}")
+    return True