Skip to content

Commit

Permalink
Added pathfinder sycl/ndpx impl
Browse files Browse the repository at this point in the history
  • Loading branch information
roxx30198 committed Nov 21, 2023
1 parent 6764112 commit 53de4a8
Show file tree
Hide file tree
Showing 10 changed files with 603 additions and 0 deletions.
5 changes: 5 additions & 0 deletions dpbench/benchmarks/rodinia/pathfinder/CMakeLists.txt
Original file line number Diff line number Diff line change
@@ -0,0 +1,5 @@
# SPDX-FileCopyrightText: 2022 - 2023 Intel Corporation
#
# SPDX-License-Identifier: Apache-2.0

add_subdirectory(pathfinder_sycl_native_ext)
32 changes: 32 additions & 0 deletions dpbench/benchmarks/rodinia/pathfinder/__init__.py
Original file line number Diff line number Diff line change
@@ -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.
"""
30 changes: 30 additions & 0 deletions dpbench/benchmarks/rodinia/pathfinder/pathfinder_initialize.py
Original file line number Diff line number Diff line change
@@ -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),
)
176 changes: 176 additions & 0 deletions dpbench/benchmarks/rodinia/pathfinder/pathfinder_numba_dpex_k.py
Original file line number Diff line number Diff line change
@@ -0,0 +1,176 @@
# 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,
rows,
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,
rows,
t,
borderCols,
block_size,
)

k = 0
for i in gpuResult[dst]:
result[k] = i
k += 1
53 changes: 53 additions & 0 deletions dpbench/benchmarks/rodinia/pathfinder/pathfinder_python.py
Original file line number Diff line number Diff line change
@@ -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]
Original file line number Diff line number Diff line change
@@ -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}
)
Original file line number Diff line number Diff line change
@@ -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"]
Loading

0 comments on commit 53de4a8

Please sign in to comment.