Added pathfinder sycl/ndpx impl

dpbench/benchmarks/rodinia/CMakeLists.txt

@@ -3,3 +3,4 @@
 # SPDX-License-Identifier: Apache-2.0
 
 add_subdirectory(gaussian)
+add_subdirectory(pathfinder)

dpbench/benchmarks/rodinia/pathfinder/CMakeLists.txt

@@ -0,0 +1,5 @@
+# SPDX-FileCopyrightText: 2022 - 2023 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+
+add_subdirectory(pathfinder_sycl_native_ext)

dpbench/benchmarks/rodinia/pathfinder/__init__.py

@@ -0,0 +1,32 @@
+# SPDX-FileCopyrightText: 2022 - 2023 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+
+"""Pathfinder Implementation.
+
+This pathfinding algorithm finds the shortest path from the first row to the last.
+
+Input
+---------
+rows<int_64> : Indicates the number of rows.
+
+cols<int_64> : Indicates the number of cols.
+
+pyramid height<int_64> : Indicates pyramid height.
+
+block_size<int_64> : Indicates block size for parallel computation.
+
+Output
+
+--------
+
+result<array<int_64>> : Indicates the minimum distance from first row to last.
+
+Method:
+
+The elements are fed to the kernel row-wise and the minimum distance is computed based
+on the minimum weight of the neighbors above.
+This is done for all rows until last and result is returned.
+
+
+"""

dpbench/benchmarks/rodinia/pathfinder/pathfinder_initialize.py

@@ -0,0 +1,30 @@
+# SPDX-FileCopyrightText: 2022 - 2023 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+"""Pathfinder initialization."""
+LOW = 0
+HIGH = 10.0
+SEED = 9
+
+
+def initialize(rows, cols, types_dict=None):
+    """Initialize the input and output matrices for pathfinder.
+
+    Args:
+        rows: number of rows.
+        cols: number of cols.
+        types_dict: data type of operand.
+
+    Returns:
+        data: input matrix.
+        result: result matrix.
+    """
+    import numpy as np
+    import numpy.random as rnd
+
+    rnd.seed(SEED)
+
+    return (
+        rnd.randint(LOW, HIGH, (rows * cols), dtype=np.int64),
+        np.empty(cols, dtype=np.int64),
+    )

dpbench/benchmarks/rodinia/pathfinder/pathfinder_numba_dpex_k.py

@@ -0,0 +1,174 @@
+# SPDX-FileCopyrightText: 2022 - 2023 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+"""Numba-dpex implementation for Pathfinder."""
+
+import dpnp
+import numba_dpex
+
+
+@numba_dpex.func
+def IN_RANGE(x, min, max):
+    """Find if x is in range.
+
+    Args:
+        x: Element to be checked.
+        min: Range min.
+        max: Range max.
+
+    Returns:
+        true/false: If in range or not.
+    """
+    return (x) >= (min) and (x) <= (max)
+
+
+@numba_dpex.func
+def min_dpex(a, b):
+    """Find the min.
+
+    Args:
+        a: first element.
+        b: second element.
+
+    Returns:
+        t: min of two.
+    """
+    t = a if a <= b else b
+    return t
+
+
+@numba_dpex.kernel
+def _pathfinder_kernel(
+    iteration,
+    gpuWall,
+    gpuSrc,
+    gpuResults,
+    cols,
+    startStep,
+    border,
+    block_size,
+):
+    """Kernel to compute the smallest element per iteration.
+
+    Args:
+        iteration: current iteration.
+        gpuWall: Matrix elements excluding that row.
+        gpuSrc: input row.
+        gpuResults: Intermediate results.
+        cols: number of columns.
+        rows: number of rows.
+        startStep: starting point.
+        border: max/min border.
+        block_size: block size for computation.
+    """
+    prev = numba_dpex.local.array((256), dtype=dpnp.int64)
+    result = numba_dpex.local.array((256), dtype=dpnp.int64)
+
+    bx = numba_dpex.get_group_id(2)
+    tx = numba_dpex.get_local_id(2)
+
+    small_block_cols = block_size - iteration * 2
+
+    blkX = small_block_cols * bx - border
+    blkXmax = blkX + block_size - 1
+
+    xidx = blkX + tx
+
+    validXmin = -blkX if blkX < 0 else 0
+    validXmax = (
+        block_size - 1 - (blkXmax - cols + 1)
+        if blkXmax > cols - 1
+        else block_size - 1
+    )
+
+    W = tx - 1
+    E = tx + 1
+
+    W = validXmin if W < validXmin else W
+    E = validXmax if E > validXmax else E
+
+    isValid = IN_RANGE(tx, validXmin, validXmax)
+
+    if IN_RANGE(xidx, 0, cols - 1):
+        prev[tx] = gpuSrc[xidx]
+
+    numba_dpex.barrier(numba_dpex.LOCAL_MEM_FENCE)
+
+    for i in range(iteration):
+        computed = False
+        if IN_RANGE(tx, i + 1, block_size - i - 2) and isValid:
+            computed = True
+            left = prev[W]
+            up = prev[tx]
+            right = prev[E]
+            shortest = min_dpex(left, up)
+            shortest = min_dpex(shortest, right)
+            index = cols * (startStep + i) + xidx
+            result[tx] = shortest + gpuWall[index]
+
+        numba_dpex.barrier(numba_dpex.LOCAL_MEM_FENCE)
+        if i == iteration - 1:
+            break
+        if computed:
+            prev[tx] = result[tx]
+        numba_dpex.barrier(numba_dpex.LOCAL_MEM_FENCE)
+
+    if computed:
+        gpuResults[xidx] = result[tx]
+
+
+def pathfinder(data, rows, cols, pyramid_height, block_size, result):
+    """Compute smallest distance from top row to bottom.
+
+    Args:
+         data: data matrix.
+         rows: number of rows.
+         cols: number of cols.
+         pyramid_height: pyramid height.
+         block_size: block size for parallel 2d-kernel.
+         result: result of operation.
+    """
+    # create a temp list that hold first row of data as first element and empty numpy array as second element
+    borderCols = pyramid_height
+    smallBlockCol = block_size - (pyramid_height) * 2
+    blockCols = int(
+        cols / smallBlockCol + (0 if cols % smallBlockCol == 0 else 1)
+    )
+    size = rows * cols
+
+    dimBlock = numba_dpex.Range(1, 1, block_size)
+    dimGrid = numba_dpex.Range(1, 1, blockCols * block_size)
+
+    gpuResult = dpnp.zeros((2, cols), dtype=dpnp.int64)
+    gpuWall = dpnp.array((size - cols), dtype=dpnp.int64)
+
+    gpuResult[0] = data[:cols]
+    gpuWall = data[cols:]
+
+    dimBlock = numba_dpex.Range(1, 1, block_size)
+    dimGrid = numba_dpex.Range(1, 1, blockCols * block_size)
+
+    src = 1
+    dst = 0
+    for t in range(0, rows - 1, pyramid_height):
+        temp = src
+        src = dst
+        dst = temp
+
+        iteration = min(pyramid_height, rows - t - 1)
+
+        _pathfinder_kernel[numba_dpex.NdRange(dimGrid, dimBlock)](
+            iteration,
+            gpuWall,
+            gpuResult[src],
+            gpuResult[dst],
+            cols,
+            t,
+            borderCols,
+            block_size,
+        )
+
+    k = 0
+    for i in gpuResult[dst]:
+        result[k] = i
+        k += 1

dpbench/benchmarks/rodinia/pathfinder/pathfinder_python.py

@@ -0,0 +1,53 @@
+# SPDX-FileCopyrightText: 2022 - 2023 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+"""Pathfinder python serial implementation."""
+
+
+def _pathfinder_kernel(device_src, device_dest, cols, current_element):
+    """Compute shortest distance per column element.
+
+    Args:
+         device_src: src row.
+         device_dest: destination row.
+         cols: number of cols.
+         current_element: Current column element.
+    """
+    left_ind = current_element - 1 if current_element >= 1 else 0
+    right_ind = current_element + 1 if current_element < cols - 1 else cols - 1
+    up_ind = current_element
+
+    left = device_src[left_ind]
+    up = device_src[up_ind]
+    right = device_src[right_ind]
+    shortest = min(left, up, right)
+
+    device_dest[current_element] += shortest
+
+
+def pathfinder(data, rows, cols, pyramid_height, block_size, result):
+    """Compute smallest distance from top row to bottom.
+
+    Args:
+         data: data matrix.
+         rows: number of rows.
+         cols: number of cols.
+         pyramid_height: pyramid height.
+         block_size: block size for parallel 2d-kernel.
+         result: result of operation.
+    """
+    # create a temp list that hold first row of data as first element and empty numpy array as second element
+    device_dest = data[:cols]  # first row
+    device_src = [0] * cols
+
+    t = 1
+    while t < rows:
+        device_src = device_dest
+        device_dest = data[t * cols : (t + 1) * cols]
+
+        for i in range(cols):
+            _pathfinder_kernel(device_src, device_dest, cols, i)
+        t += 1
+
+    for i in range(cols):
+        result[i] = device_dest[i]

dpbench/benchmarks/rodinia/pathfinder/pathfinder_sycl_native_ext/CMakeLists.txt

@@ -0,0 +1,14 @@
+# SPDX-FileCopyrightText: 2022 - 2023 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+
+set(module_name pathfinder_sycl)
+set(py_module_name _${module_name})
+python_add_library(${py_module_name} MODULE ${module_name}/${py_module_name}.cpp)
+add_sycl_to_target(TARGET ${py_module_name} SOURCES ${module_name}/${py_module_name}.cpp)
+target_include_directories(${py_module_name} PRIVATE ${Dpctl_INCLUDE_DIRS})
+
+file(RELATIVE_PATH py_module_dest ${CMAKE_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR})
+install(TARGETS ${py_module_name}
+  DESTINATION ${py_module_dest}/${module_name}
+)

dpbench/benchmarks/rodinia/pathfinder/pathfinder_sycl_native_ext/__init__.py

@@ -0,0 +1,8 @@
+# SPDX-FileCopyrightText: 2022 - 2023 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+"""Sycl implementation for Pathfinder."""
+
+from .pathfinder_sycl._pathfinder_sycl import pathfinder as pathfinder_sycl
+
+__all__ = ["pathfinder_sycl"]