-
Notifications
You must be signed in to change notification settings - Fork 33
Commit
This commit does not belong to any branch on this repository, and may belong to a fork outside of the repository.
Merge pull request #888 from chudur-budur/github-871
A more consistent kernel launch parameter syntax
- Loading branch information
Showing
28 changed files
with
427 additions
and
106 deletions.
There are no files selected for viewing
This file contains 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 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 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 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.