-
Notifications
You must be signed in to change notification settings - Fork 32
A more consistent kernel launch parameter syntax #888
New issue
Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.
By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.
Already on GitHub? Sign in to your account
Merged
Merged
Changes from all commits
Commits
Show all changes
6 commits
Select commit
Hold shift + click to select a range
3582e23
Add wrapper classes for Range and NdRange.
1e72a21
Fix kernel lauch parameter syntax
chudur-budur d0e46f2
Fixed kernel lauch params and added NdRange in Jitkernel
chudur-budur e66c0fc
Getting rid of core.kernel_interface.indexers.NdRange, using core.ker…
chudur-budur 64b81a9
A better implementation for the Range/NdRange class, with examples
chudur-budur 57c7d10
Updating all examples to work with Range/NdRange
chudur-budur File filter
Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
There are no files selected for viewing
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,218 @@ | ||
from collections.abc import Iterable | ||
|
||
|
||
class Range(tuple): | ||
"""A data structure to encapsulate a single kernel lauch parameter. | ||
|
||
The range is an abstraction that describes the number of elements | ||
in each dimension of buffers and index spaces. It can contain | ||
1, 2, or 3 numbers, dependending on the dimensionality of the | ||
object it describes. | ||
|
||
This is just a wrapper class on top of a 3-tuple. The kernel launch | ||
parameter is consisted of three int's. This class basically mimics | ||
the behavior of `sycl::range`. | ||
""" | ||
|
||
def __new__(cls, dim0, dim1=None, dim2=None): | ||
"""Constructs a 1, 2, or 3 dimensional range. | ||
|
||
Args: | ||
dim0 (int): The range of the first dimension. | ||
dim1 (int, optional): The range of second dimension. | ||
Defaults to None. | ||
dim2 (int, optional): The range of the third dimension. | ||
Defaults to None. | ||
|
||
Raises: | ||
TypeError: If dim0 is not an int. | ||
TypeError: If dim1 is not an int. | ||
TypeError: If dim2 is not an int. | ||
""" | ||
if not isinstance(dim0, int): | ||
raise TypeError("dim0 of a Range must be an int.") | ||
_values = [dim0] | ||
if dim1: | ||
if not isinstance(dim1, int): | ||
raise TypeError("dim1 of a Range must be an int.") | ||
_values.append(dim1) | ||
if dim2: | ||
if not isinstance(dim2, int): | ||
raise TypeError("dim2 of a Range must be an int.") | ||
_values.append(dim2) | ||
return super(Range, cls).__new__(cls, tuple(_values)) | ||
|
||
def get(self, index): | ||
"""Returns the range of a single dimension. | ||
|
||
Args: | ||
index (int): The index of the dimension, i.e. [0,2] | ||
|
||
Returns: | ||
int: The range of the dimension indexed by `index`. | ||
""" | ||
return self[index] | ||
|
||
def size(self): | ||
"""Returns the size of a range. | ||
|
||
Returns the size of a range by multiplying | ||
the range of the individual dimensions. | ||
|
||
Returns: | ||
int: The size of a range. | ||
""" | ||
n = len(self) | ||
if n > 2: | ||
return self[0] * self[1] * self[2] | ||
elif n > 1: | ||
return self[0] * self[1] | ||
else: | ||
return self[0] | ||
|
||
|
||
class NdRange: | ||
"""A class to encapsulate all kernel launch parameters. | ||
|
||
The NdRange defines the index space for a work group as well as | ||
the global index space. It is passed to parallel_for to execute | ||
a kernel on a set of work items. | ||
|
||
This class basically contains two Range object, one for the global_range | ||
and the other for the local_range. The global_range parameter contains | ||
the global index space and the local_range parameter contains the index | ||
space of a work group. This class mimics the behavior of `sycl::nd_range` | ||
class. | ||
""" | ||
|
||
def __init__(self, global_size, local_size): | ||
"""Constructor for NdRange class. | ||
|
||
Args: | ||
global_size (Range or tuple of int's): The values for | ||
the global_range. | ||
local_size (Range or tuple of int's, optional): The values for | ||
the local_range. Defaults to None. | ||
""" | ||
if isinstance(global_size, Range): | ||
self._global_range = global_size | ||
elif isinstance(global_size, Iterable): | ||
self._global_range = Range(*global_size) | ||
else: | ||
TypeError("Unknwon argument type for NdRange global_size.") | ||
|
||
if isinstance(local_size, Range): | ||
self._local_range = local_size | ||
elif isinstance(local_size, Iterable): | ||
self._local_range = Range(*local_size) | ||
else: | ||
TypeError("Unknwon argument type for NdRange local_size.") | ||
|
||
@property | ||
def global_range(self): | ||
"""Accessor for global_range. | ||
|
||
Returns: | ||
Range: The `global_range` `Range` object. | ||
""" | ||
return self._global_range | ||
|
||
@property | ||
def local_range(self): | ||
"""Accessor for local_range. | ||
|
||
Returns: | ||
Range: The `local_range` `Range` object. | ||
""" | ||
return self._local_range | ||
|
||
def get_global_range(self): | ||
"""Returns a Range defining the index space. | ||
|
||
Returns: | ||
Range: A `Range` object defining the index space. | ||
""" | ||
return self._global_range | ||
|
||
def get_local_range(self): | ||
"""Returns a Range defining the index space of a work group. | ||
|
||
Returns: | ||
Range: A `Range` object to specify index space of a work group. | ||
""" | ||
return self._local_range | ||
|
||
def __str__(self): | ||
"""str() function for NdRange class. | ||
|
||
Returns: | ||
str: str representation for NdRange class. | ||
""" | ||
return ( | ||
"(" + str(self._global_range) + ", " + str(self._local_range) + ")" | ||
) | ||
|
||
def __repr__(self): | ||
"""repr() function for NdRange class. | ||
|
||
Returns: | ||
str: str representation for NdRange class. | ||
""" | ||
return self.__str__() | ||
|
||
|
||
if __name__ == "__main__": | ||
r1 = Range(1) | ||
print("r1 =", r1) | ||
|
||
r2 = Range(1, 2) | ||
print("r2 =", r2) | ||
|
||
r3 = Range(1, 2, 3) | ||
print("r3 =", r3, ", len(r3) =", len(r3)) | ||
|
||
r3 = Range(*(1, 2, 3)) | ||
print("r3 =", r3, ", len(r3) =", len(r3)) | ||
|
||
r3 = Range(*[1, 2, 3]) | ||
print("r3 =", r3, ", len(r3) =", len(r3)) | ||
|
||
print("r1.get(0) =", r1.get(0)) | ||
try: | ||
print("r2.get(2) =", r2.get(2)) | ||
except Exception as e: | ||
print(e) | ||
|
||
print("r3.get(0) =", r3.get(0)) | ||
print("r3.get(1) =", r3.get(1)) | ||
|
||
print("r1[0] =", r1[0]) | ||
try: | ||
print("r2[2] =", r2[2]) | ||
except Exception as e: | ||
print(e) | ||
|
||
print("r3[0] =", r3[0]) | ||
print("r3[1] =", r3[1]) | ||
|
||
try: | ||
r4 = Range(1, 2, 3, 4) | ||
except Exception as e: | ||
print(e) | ||
|
||
try: | ||
r5 = Range(*(1, 2, 3, 4)) | ||
except Exception as e: | ||
print(e) | ||
|
||
ndr1 = NdRange(Range(1, 2)) | ||
print("ndr1 =", ndr1) | ||
|
||
ndr2 = NdRange(Range(1, 2), Range(1, 1, 1)) | ||
print("ndr2 =", ndr2) | ||
|
||
ndr3 = NdRange((1, 2)) | ||
print("ndr3 =", ndr3) | ||
|
||
ndr4 = NdRange((1, 2), (1, 1, 1)) | ||
print("ndr4 =", ndr4) |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Oops, something went wrong.
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
Uh oh!
There was an error while loading. Please reload this page.