From 4ab2e9f247c0c511da42d96808c1ae3dd03c9aea Mon Sep 17 00:00:00 2001 From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com> Date: Fri, 18 Nov 2022 13:02:00 +0100 Subject: [PATCH 01/29] Add cudaMemcpyKind enum and fix typo in printing --- pyccel/codegen/printing/ccudacode.py | 2 +- pyccel/stdlib/cuda_ndarrays/cuda_ndarrays.h | 8 ++++++++ 2 files changed, 9 insertions(+), 1 deletion(-) diff --git a/pyccel/codegen/printing/ccudacode.py b/pyccel/codegen/printing/ccudacode.py index d9c83f6458..b7a366481a 100644 --- a/pyccel/codegen/printing/ccudacode.py +++ b/pyccel/codegen/printing/ccudacode.py @@ -524,7 +524,7 @@ def copy_CudaArray_Data(self, expr): symbol=expr, severity='error') memcpy_kind_src = str(rhs.current_location).capitalize() memcpy_kind_dest = 'Host' if rhs.memory_location == 'host' else 'Device' - memcpy_kind = "{}To{}".format(memcpy_kind_src, memcpy_kind_dest) + memcpy_kind = "cudaMemcpy{}To{}".format(memcpy_kind_src, memcpy_kind_dest) if rhs.rank > 1: # flattening the args to use them in C initialization. arg = self._flatten_list(arg) diff --git a/pyccel/stdlib/cuda_ndarrays/cuda_ndarrays.h b/pyccel/stdlib/cuda_ndarrays/cuda_ndarrays.h index d0cce3b5f0..14ef818a72 100644 --- a/pyccel/stdlib/cuda_ndarrays/cuda_ndarrays.h +++ b/pyccel/stdlib/cuda_ndarrays/cuda_ndarrays.h @@ -3,6 +3,14 @@ #include "../ndarrays/ndarrays.h" +enum e_cudaMemcpyKind { + cudaMemcpyHostToHost = 0, + cudaMemcpyHostToDevice = 1, + cudaMemcpyDeviceToHost = 2, + cudaMemcpyDeviceToDevice = 3, + cudaMemcpyDefault = 4 +}; + __global__ void cuda_array_arange_int8(t_ndarray arr, int start); __global__ From 69cb9d8c97ba90a05a5689fc8b8bf6cf50e59704 Mon Sep 17 00:00:00 2001 From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com> Date: Fri, 18 Nov 2022 20:08:48 +0100 Subject: [PATCH 02/29] Addition of context in array creation --- pyccel/ast/cudaext.py | 11 +++++++++-- pyccel/ast/cupyext.py | 17 ++++++++++++++--- pyccel/ast/variable.py | 15 +++++++++++++-- pyccel/codegen/printing/ccudacode.py | 8 ++++---- pyccel/parser/semantic.py | 27 +++++++++++++++++++++++---- 5 files changed, 63 insertions(+), 15 deletions(-) diff --git a/pyccel/ast/cudaext.py b/pyccel/ast/cudaext.py index 79f42bda74..7b58ba6e1b 100644 --- a/pyccel/ast/cudaext.py +++ b/pyccel/ast/cudaext.py @@ -74,11 +74,11 @@ class CudaArray(CudaNewArray): arg : list, tuple, PythonList """ - __slots__ = ('_arg','_dtype','_precision','_shape','_rank','_order','_memory_location') + __slots__ = ('_arg','_dtype','_precision','_shape','_rank','_order','_memory_location', '_current_context') _attribute_nodes = ('_arg',) name = 'array' - def __init__(self, arg, dtype=None, order='C', memory_location='managed'): + def __init__(self, arg, dtype=None, order='C', current_context='host', memory_location='managed'): if not isinstance(arg, (PythonTuple, PythonList, Variable)): raise TypeError('Unknown type of %s.' % type(arg)) @@ -118,6 +118,9 @@ def __init__(self, arg, dtype=None, order='C', memory_location='managed'): #Verify memory location if memory_location not in ('host', 'device', 'managed'): raise ValueError("memory_location must be 'host', 'device' or 'managed'") + if current_context not in ('host', 'device'): + raise ValueError("The current context can only be 'host' or 'device'") + self._arg = arg self._shape = shape self._rank = rank @@ -125,6 +128,7 @@ def __init__(self, arg, dtype=None, order='C', memory_location='managed'): self._order = order self._precision = prec self._memory_location = memory_location + self._current_context = current_context super().__init__() def __str__(self): @@ -136,6 +140,9 @@ def arg(self): @property def memory_location(self): return self._memory_location + @property + def current_context(self): + return self._current_context class CudaDeviceSynchronize(PyccelInternalFunction): "Represents a call to Cuda.deviceSynchronize for code generation." diff --git a/pyccel/ast/cupyext.py b/pyccel/ast/cupyext.py index 08650db3eb..30144c39ea 100644 --- a/pyccel/ast/cupyext.py +++ b/pyccel/ast/cupyext.py @@ -65,11 +65,11 @@ class CupyArray(CudaNewArray): arg : list, tuple, PythonList """ - __slots__ = ('_arg','_dtype','_precision','_shape','_rank','_order') + __slots__ = ('_arg', '_dtype','_precision','_shape','_rank','_order', '_memory_location', '_current_context') _attribute_nodes = ('_arg',) name = 'array' - def __init__(self, arg, dtype=None, order='C'): + def __init__(self, arg, dtype=None, order='C', current_context='host'): if not isinstance(arg, (PythonTuple, PythonList, Variable)): raise TypeError('Unknown type of %s.' % type(arg)) @@ -92,7 +92,8 @@ def __init__(self, arg, dtype=None, order='C'): shape = process_shape(False, arg.shape) rank = len(shape) - + if current_context not in ('host', 'device'): + raise ValueError("The current context can only be 'host' or 'device'") if rank < 2: order = None else: @@ -113,6 +114,8 @@ def __init__(self, arg, dtype=None, order='C'): self._dtype = dtype self._order = order self._precision = prec + self._current_context = current_context + self._memory_location = 'device' super().__init__() def __str__(self): @@ -122,6 +125,14 @@ def __str__(self): def arg(self): return self._arg + @property + def current_context(self): + return self._current_context + + @property + def memory_location(self): + return self._memory_location + #============================================================================== class CupyArange(CudaNewArray): """ diff --git a/pyccel/ast/variable.py b/pyccel/ast/variable.py index 36a92ac858..ccd2c88cff 100644 --- a/pyccel/ast/variable.py +++ b/pyccel/ast/variable.py @@ -106,8 +106,8 @@ class base if variable is an object or an object member [Default value: None] >>> Variable('int', DottedName('matrix', 'n_rows')) matrix.n_rows """ - __slots__ = ('_name', '_alloc_shape', '_memory_handling', '_memory_location', '_is_const', - '_is_target', '_is_optional', '_allows_negative_indexes', + __slots__ = ('_name', '_alloc_shape', '_memory_handling', '_memory_location', '_current_context', + '_is_const', '_is_target', '_is_optional', '_allows_negative_indexes', '_cls_base', '_is_argument', '_is_kwonly', '_is_temp','_dtype','_precision', '_rank','_shape','_order','_is_private') _attribute_nodes = () @@ -120,6 +120,7 @@ def __init__( rank=0, memory_handling='stack', memory_location='host', + current_context = 'host', is_const=False, is_target=False, is_optional=False, @@ -159,6 +160,10 @@ def __init__( raise ValueError("memory_location must be 'host', 'device' or 'managed'") self._memory_location = memory_location + if current_context not in ('host', 'device'): + raise ValueError("The current context can only be 'host' or 'device'") + self._current_context = current_context + if not isinstance(is_const, bool): raise TypeError('is_const must be a boolean.') self._is_const = is_const @@ -320,6 +325,12 @@ def memory_location(self): """ return self._memory_location + @property + def current_context(self): + """ Indicates if the variable is currently in the host or device context + """ + return self._current_context + @memory_location.setter def memory_location(self, memory_location): if memory_location not in ('host', 'device', 'managed'): diff --git a/pyccel/codegen/printing/ccudacode.py b/pyccel/codegen/printing/ccudacode.py index 73df5043f2..a154760fa0 100644 --- a/pyccel/codegen/printing/ccudacode.py +++ b/pyccel/codegen/printing/ccudacode.py @@ -520,10 +520,10 @@ def copy_CudaArray_Data(self, expr): dtype = self.find_in_ndarray_type_registry(self._print(rhs.dtype), rhs.precision) arg = rhs.arg if isinstance(rhs, (CudaArray, CupyArray)) else rhs - if rhs.current_location == 'device' and rhs.memory_location == 'host': + if rhs.current_context == 'device' and rhs.memory_location == 'host': errors.report("You can't create a host array from a device function", symbol=expr, severity='error') - memcpy_kind_src = str(rhs.current_location).capitalize() + memcpy_kind_src = str(rhs.current_context).capitalize() memcpy_kind_dest = 'Host' if rhs.memory_location == 'host' else 'Device' memcpy_kind = "cudaMemcpy{}To{}".format(memcpy_kind_src, memcpy_kind_dest) @@ -534,12 +534,12 @@ def copy_CudaArray_Data(self, expr): self.add_import(c_imports['string']) if isinstance(arg, Variable): arg = self._print(arg) - cpy_data = "cudaMemcpy({0}.raw_data, {1}.{2}, {0}.buffer_size, cudaMemcpyHostToDevice);".format(lhs, arg, dtype) + cpy_data = "cudaMemcpy({0}.raw_data, {1}.{2}, {0}.buffer_size, {3});".format(lhs, arg, dtype, memcpy_kind) return '%s\n' % (cpy_data) else : arg = ', '.join(self._print(i) for i in arg) dummy_array = "%s %s[] = {%s};\n" % (declare_dtype, dummy_array_name, arg) - cpy_data = "cudaMemcpy({0}.raw_data, {1}, {0}.buffer_size, cudaMemcpyHostToDevice);".format(self._print(lhs), dummy_array_name, dtype) + cpy_data = "cudaMemcpy({0}.raw_data, {1}, {0}.buffer_size, {3});".format(self._print(lhs), dummy_array_name, dtype, memcpy_kind) return '%s%s\n' % (dummy_array, cpy_data) def _print_CudaDeviceSynchronize(self, expr): diff --git a/pyccel/parser/semantic.py b/pyccel/parser/semantic.py index 53c022993f..92f4c4c518 100644 --- a/pyccel/parser/semantic.py +++ b/pyccel/parser/semantic.py @@ -61,6 +61,7 @@ from pyccel.ast.core import Assert from pyccel.ast.class_defs import NumpyArrayClass, TupleClass, get_cls_base, CudaArrayClass +from pyccel.ast.cupyext import CupyArray from pyccel.ast.datatypes import NativeRange, str_dtype from pyccel.ast.datatypes import NativeSymbol @@ -489,11 +490,29 @@ def _infere_type(self, expr, **settings): d_var['precision' ] = expr.precision d_var['cls_base' ] = NumpyArrayClass return d_var - + elif isinstance(expr, CupyArray): + d_var['datatype' ] = expr.dtype + d_var['memory_handling'] = 'heap' if expr.rank > 0 else 'stack' + if 'device' in self.scope.decorators or\ + 'kernel' in self.scope.decorators: + d_var['current_context'] = 'device' + else: + d_var['current_context'] = 'device' + d_var['shape' ] = expr.shape + d_var['rank' ] = expr.rank + d_var['order' ] = expr.order + d_var['precision' ] = expr.precision + d_var['cls_base' ] = CudaArrayClass + return d_var elif isinstance(expr, CudaNewArray): d_var['datatype' ] = expr.dtype d_var['memory_handling'] = 'heap' if expr.rank > 0 else 'stack' d_var['memory_location'] = expr.memory_location + if 'device' in self.scope.decorators or\ + 'kernel' in self.scope.decorators: + d_var['current_context'] = 'device' + else: + d_var['current_context'] = 'device' d_var['shape' ] = expr.shape d_var['rank' ] = expr.rank d_var['order' ] = expr.order @@ -812,7 +831,7 @@ def _handle_function(self, expr, func, args, **settings): func = func.cls_name if func in (CudaThreadIdx, CudaBlockDim, CudaBlockIdx, CudaGridDim): if 'kernel' not in self.scope.decorators\ - or 'device' not in self.scope.decorators: + and 'device' not in self.scope.decorators: errors.report("Cuda internal variables should only be used in Kernel or Device functions", symbol = expr, severity = 'fatal') @@ -902,11 +921,11 @@ def _handle_kernel(self, expr, func, args, **settings): symbol = expr, severity='fatal') # TODO : type check the NUMBER OF BLOCKS 'numBlocks' and threads per block 'tpblock' - if not isinstance(expr.numBlocks, LiteralInteger): + if not isinstance(expr.numBlocks, (LiteralInteger, PyccelSymbol)): errors.report("Invalid Block number parameter for Kernel call", symbol = expr, severity='error') - if not isinstance(expr.tpblock, LiteralInteger): + if not isinstance(expr.tpblock, (LiteralInteger, PyccelSymbol)): errors.report("Invalid Thread per Block parameter for Kernel call", symbol = expr, severity='error') From 77b25a0be0037d13818a7759737a562def16536c Mon Sep 17 00:00:00 2001 From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com> Date: Fri, 18 Nov 2022 22:36:12 +0100 Subject: [PATCH 03/29] Choose context using decorators --- pyccel/parser/semantic.py | 10 +++++++++- 1 file changed, 9 insertions(+), 1 deletion(-) diff --git a/pyccel/parser/semantic.py b/pyccel/parser/semantic.py index 92f4c4c518..725140a562 100644 --- a/pyccel/parser/semantic.py +++ b/pyccel/parser/semantic.py @@ -95,7 +95,7 @@ from pyccel.ast.numpyext import NumpyNewArray, NumpyNonZero from pyccel.ast.numpyext import DtypePrecisionToCastFunction -from pyccel.ast.cudaext import CudaNewArray, CudaThreadIdx, CudaBlockDim, CudaBlockIdx, CudaGridDim +from pyccel.ast.cudaext import CudaArray, CudaNewArray, CudaThreadIdx, CudaBlockDim, CudaBlockIdx, CudaGridDim, CudaToDevice from pyccel.ast.omp import (OMP_For_Loop, OMP_Simd_Construct, OMP_Distribute_Construct, OMP_TaskLoop_Construct, OMP_Sections_Construct, Omp_End_Clause, @@ -836,6 +836,14 @@ def _handle_function(self, expr, func, args, **settings): symbol = expr, severity = 'fatal') + if func in (CupyArray, CudaArray): + if not 'current_context' in args: + if 'device' in self.scope.decorators or\ + 'kernel' in self.scope.decorators: + new_arg = FunctionCallArgument('device', 'current_context',) + else: + new_arg = FunctionCallArgument('host', 'current_context',) + args.append(new_arg) args, kwargs = split_positional_keyword_arguments(*args) for a in args: if getattr(a,'dtype',None) == 'tuple': From 521d19908b8c83179ead0d48fe84c097566fecb0 Mon Sep 17 00:00:00 2001 From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com> Date: Sat, 19 Nov 2022 02:01:41 +0100 Subject: [PATCH 04/29] Add cuda to_host and to_device --- pyccel/ast/cudaext.py | 18 ++++++++++++++++++ pyccel/codegen/printing/ccudacode.py | 3 --- pyccel/parser/semantic.py | 25 ++++--------------------- 3 files changed, 22 insertions(+), 24 deletions(-) diff --git a/pyccel/ast/cudaext.py b/pyccel/ast/cudaext.py index 7b58ba6e1b..36c55929d4 100644 --- a/pyccel/ast/cudaext.py +++ b/pyccel/ast/cudaext.py @@ -118,6 +118,7 @@ def __init__(self, arg, dtype=None, order='C', current_context='host', memory_lo #Verify memory location if memory_location not in ('host', 'device', 'managed'): raise ValueError("memory_location must be 'host', 'device' or 'managed'") + # ... if current_context not in ('host', 'device'): raise ValueError("The current context can only be 'host' or 'device'") @@ -197,10 +198,27 @@ def __new__(cls, dim=0): return expr[0] return PythonTuple(*expr) +class CudaToDevice(CudaArray): + def __init__(self, arg): + if not isinstance(arg, Variable) or not arg.is_ndarray: + raise TypeError("Argument must be an ndarray variable") + order = arg.order + current_context = 'device' if not arg.memory_location == 'host' else arg.memory_location + super().__init__(arg, order=order, current_context=current_context, memory_location='device') + +class CudaToHost(CudaArray): + def __init__(self, arg): + if not isinstance(arg, Variable) or not arg.is_ndarray: + raise TypeError("Argument must be an ndarray variable") + order = arg.order + current_context = 'device' if not arg.memory_location == 'host' else arg.memory_location + super().__init__(arg, order=order, current_context=current_context, memory_location='host') cuda_funcs = { # 'deviceSynchronize' : CudaDeviceSynchronize, 'array' : PyccelFunctionDef('array' , CudaArray), + 'to_device' : PyccelFunctionDef('to_device' , CudaToDevice), + 'to_host' : PyccelFunctionDef('to_host' , CudaToHost), 'deviceSynchronize' : PyccelFunctionDef('deviceSynchronize' , CudaDeviceSynchronize), 'threadIdx' : PyccelFunctionDef('threadIdx' , CudaThreadIdx), 'blockDim' : PyccelFunctionDef('blockDim' , CudaBlockDim), diff --git a/pyccel/codegen/printing/ccudacode.py b/pyccel/codegen/printing/ccudacode.py index a154760fa0..3b794eb3aa 100644 --- a/pyccel/codegen/printing/ccudacode.py +++ b/pyccel/codegen/printing/ccudacode.py @@ -520,9 +520,6 @@ def copy_CudaArray_Data(self, expr): dtype = self.find_in_ndarray_type_registry(self._print(rhs.dtype), rhs.precision) arg = rhs.arg if isinstance(rhs, (CudaArray, CupyArray)) else rhs - if rhs.current_context == 'device' and rhs.memory_location == 'host': - errors.report("You can't create a host array from a device function", - symbol=expr, severity='error') memcpy_kind_src = str(rhs.current_context).capitalize() memcpy_kind_dest = 'Host' if rhs.memory_location == 'host' else 'Device' memcpy_kind = "cudaMemcpy{}To{}".format(memcpy_kind_src, memcpy_kind_dest) diff --git a/pyccel/parser/semantic.py b/pyccel/parser/semantic.py index 725140a562..ab65b391c4 100644 --- a/pyccel/parser/semantic.py +++ b/pyccel/parser/semantic.py @@ -490,29 +490,12 @@ def _infere_type(self, expr, **settings): d_var['precision' ] = expr.precision d_var['cls_base' ] = NumpyArrayClass return d_var - elif isinstance(expr, CupyArray): - d_var['datatype' ] = expr.dtype - d_var['memory_handling'] = 'heap' if expr.rank > 0 else 'stack' - if 'device' in self.scope.decorators or\ - 'kernel' in self.scope.decorators: - d_var['current_context'] = 'device' - else: - d_var['current_context'] = 'device' - d_var['shape' ] = expr.shape - d_var['rank' ] = expr.rank - d_var['order' ] = expr.order - d_var['precision' ] = expr.precision - d_var['cls_base' ] = CudaArrayClass - return d_var + elif isinstance(expr, CudaNewArray): d_var['datatype' ] = expr.dtype d_var['memory_handling'] = 'heap' if expr.rank > 0 else 'stack' d_var['memory_location'] = expr.memory_location - if 'device' in self.scope.decorators or\ - 'kernel' in self.scope.decorators: - d_var['current_context'] = 'device' - else: - d_var['current_context'] = 'device' + d_var['current_context'] = expr.current_context d_var['shape' ] = expr.shape d_var['rank' ] = expr.rank d_var['order' ] = expr.order @@ -840,9 +823,9 @@ def _handle_function(self, expr, func, args, **settings): if not 'current_context' in args: if 'device' in self.scope.decorators or\ 'kernel' in self.scope.decorators: - new_arg = FunctionCallArgument('device', 'current_context',) + new_arg = FunctionCallArgument('device', 'current_context') else: - new_arg = FunctionCallArgument('host', 'current_context',) + new_arg = FunctionCallArgument('host', 'current_context') args.append(new_arg) args, kwargs = split_positional_keyword_arguments(*args) for a in args: From ca9c7d1ea7f325f1bdeb6fb8d387a4a8e5435d82 Mon Sep 17 00:00:00 2001 From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com> Date: Sat, 19 Nov 2022 02:47:13 +0100 Subject: [PATCH 05/29] Raise error when creating host array from kernel --- pyccel/parser/semantic.py | 21 +++++++++++++-------- 1 file changed, 13 insertions(+), 8 deletions(-) diff --git a/pyccel/parser/semantic.py b/pyccel/parser/semantic.py index ab65b391c4..b86baa65d6 100644 --- a/pyccel/parser/semantic.py +++ b/pyccel/parser/semantic.py @@ -819,15 +819,20 @@ def _handle_function(self, expr, func, args, **settings): symbol = expr, severity = 'fatal') - if func in (CupyArray, CudaArray): - if not 'current_context' in args: - if 'device' in self.scope.decorators or\ - 'kernel' in self.scope.decorators: - new_arg = FunctionCallArgument('device', 'current_context') - else: - new_arg = FunctionCallArgument('host', 'current_context') - args.append(new_arg) args, kwargs = split_positional_keyword_arguments(*args) + if func in (CupyArray, CudaArray, NumpyArray): + if 'device' in self.scope.decorators or 'kernel' in self.scope.decorators: + current_context = 'device' + if ('memory_location' in kwargs.keys() and kwargs['memory_location'] == 'host')\ + or func is NumpyArray: + errors.report("Host arrays cannot be allocated on the Device", + symbol = expr, + severity = 'fatal') + else: + current_context = 'host' + if func in (CudaArray, CupyArray): + kwargs['current_context'] = current_context + for a in args: if getattr(a,'dtype',None) == 'tuple': self._infere_type(a, **settings) From 938c732b46ab83d5a4bcbd7c11e940ced369f5cc Mon Sep 17 00:00:00 2001 From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com> Date: Sat, 19 Nov 2022 03:53:53 +0100 Subject: [PATCH 06/29] Create temporary variables when using numpy/cuda/cupy array as arg --- pyccel/parser/semantic.py | 10 ++++++++-- 1 file changed, 8 insertions(+), 2 deletions(-) diff --git a/pyccel/parser/semantic.py b/pyccel/parser/semantic.py index b86baa65d6..3b609b7ff7 100644 --- a/pyccel/parser/semantic.py +++ b/pyccel/parser/semantic.py @@ -95,13 +95,13 @@ from pyccel.ast.numpyext import NumpyNewArray, NumpyNonZero from pyccel.ast.numpyext import DtypePrecisionToCastFunction -from pyccel.ast.cudaext import CudaArray, CudaNewArray, CudaThreadIdx, CudaBlockDim, CudaBlockIdx, CudaGridDim, CudaToDevice +from pyccel.ast.cudaext import CudaArray, CudaNewArray, CudaThreadIdx, CudaBlockDim, CudaBlockIdx, CudaGridDim from pyccel.ast.omp import (OMP_For_Loop, OMP_Simd_Construct, OMP_Distribute_Construct, OMP_TaskLoop_Construct, OMP_Sections_Construct, Omp_End_Clause, OMP_Single_Construct) -from pyccel.ast.operators import PyccelIs, PyccelIsNot, IfTernaryOperator, PyccelUnarySub +from pyccel.ast.operators import PyccelArithmeticOperator, PyccelIs, PyccelIsNot, IfTernaryOperator, PyccelUnarySub from pyccel.ast.operators import PyccelNot, PyccelEq, PyccelAdd, PyccelMul, PyccelPow from pyccel.ast.operators import PyccelAssociativeParenthesis, PyccelDiv @@ -721,6 +721,12 @@ def _handle_function_args(self, arguments, **settings): if isinstance(a.value, StarredArguments): args.extend([FunctionCallArgument(av) for av in a.value.args_var]) else: + if isinstance(a.value, PyccelArithmeticOperator) and a.value.rank\ + or isinstance(a.value, (NumpyNewArray, CudaNewArray)): + tmp_var = PyccelSymbol(self.scope.get_new_name(), is_temp=True) + assign = self._visit(Assign(tmp_var, arg.value, fst= arg.value.fst)) + self._additional_exprs[-1].append(assign) + a = FunctionCallArgument(self._visit(tmp_var)) args.append(a) return args From 423571a20a35402466ea29d298f9907d94ae7916 Mon Sep 17 00:00:00 2001 From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com> Date: Sat, 19 Nov 2022 04:56:17 +0100 Subject: [PATCH 07/29] Add docstring to CudaToHost, CudaToDevice and CudaGrid --- pyccel/ast/cudaext.py | 11 +++++++++++ 1 file changed, 11 insertions(+) diff --git a/pyccel/ast/cudaext.py b/pyccel/ast/cudaext.py index 36c55929d4..53dd4c69c5 100644 --- a/pyccel/ast/cudaext.py +++ b/pyccel/ast/cudaext.py @@ -187,6 +187,11 @@ class CudaBlockDim(CudaInternalVar) : pass class CudaBlockIdx(CudaInternalVar) : pass class CudaGridDim(CudaInternalVar) : pass class CudaGrid(PyccelAstNode) : + """Represents a call to cuda.grid for code generation. + dim : int + The number of dimensions of requested, it should correspond + to the number of dimensions declared when instantiating the kernel + """ def __new__(cls, dim=0): if not isinstance(dim, LiteralInteger): raise TypeError("dimension need to be an integer") @@ -199,6 +204,9 @@ def __new__(cls, dim=0): return PythonTuple(*expr) class CudaToDevice(CudaArray): + """Represents a call to cuda.to_device for code generation. + arg : Variable + """ def __init__(self, arg): if not isinstance(arg, Variable) or not arg.is_ndarray: raise TypeError("Argument must be an ndarray variable") @@ -208,6 +216,9 @@ def __init__(self, arg): class CudaToHost(CudaArray): def __init__(self, arg): + """Represents a call to cuda.to_host for code generation. + arg : Variable + """ if not isinstance(arg, Variable) or not arg.is_ndarray: raise TypeError("Argument must be an ndarray variable") order = arg.order From 7a703ae52bc2536ac19d1cbe35c62a36d04de77f Mon Sep 17 00:00:00 2001 From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com> Date: Sat, 19 Nov 2022 05:10:15 +0100 Subject: [PATCH 08/29] Add test for cuda.to_device and cuda.to_host --- .../scripts/ccuda/cuda_to_host_device.py | 23 +++++++++++++++++++ 1 file changed, 23 insertions(+) create mode 100644 tests/internal/scripts/ccuda/cuda_to_host_device.py diff --git a/tests/internal/scripts/ccuda/cuda_to_host_device.py b/tests/internal/scripts/ccuda/cuda_to_host_device.py new file mode 100644 index 0000000000..2f55c46e7a --- /dev/null +++ b/tests/internal/scripts/ccuda/cuda_to_host_device.py @@ -0,0 +1,23 @@ +from pyccel.decorators import kernel, types +from pyccel import cuda +import numpy as np +import cupy + +@kernel +@types('int[:]', 'int[:]') +def mult(a, b): + index = cuda.blockIdx(0) * cuda.blockDim(0) + cuda.threadIdx(0) + a[index] = b[index] * a[index] + +if __name__ == '__main__': + threads_per_block = 5 + n_blocks = 1 + a = np.array([0,1,2,3,4]) + b = cuda.to_device(a) + c = cuda.to_device(np.array([4,3,2,1,0])) + cuda.deviceSynchronize() + mult[n_blocks, threads_per_block](a, c) + cuda.deviceSynchronize() + c = cuda.to_host(b) + print(c) + \ No newline at end of file From f4d773b82e9675cd6e96c2beb53f31814dfecbe9 Mon Sep 17 00:00:00 2001 From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com> Date: Sat, 3 Dec 2022 23:13:24 +0100 Subject: [PATCH 09/29] Remove numba functions from cudaext --- pyccel/ast/cudaext.py | 24 ------------------------ 1 file changed, 24 deletions(-) diff --git a/pyccel/ast/cudaext.py b/pyccel/ast/cudaext.py index 53dd4c69c5..45f00d8e81 100644 --- a/pyccel/ast/cudaext.py +++ b/pyccel/ast/cudaext.py @@ -203,33 +203,9 @@ def __new__(cls, dim=0): return expr[0] return PythonTuple(*expr) -class CudaToDevice(CudaArray): - """Represents a call to cuda.to_device for code generation. - arg : Variable - """ - def __init__(self, arg): - if not isinstance(arg, Variable) or not arg.is_ndarray: - raise TypeError("Argument must be an ndarray variable") - order = arg.order - current_context = 'device' if not arg.memory_location == 'host' else arg.memory_location - super().__init__(arg, order=order, current_context=current_context, memory_location='device') - -class CudaToHost(CudaArray): - def __init__(self, arg): - """Represents a call to cuda.to_host for code generation. - arg : Variable - """ - if not isinstance(arg, Variable) or not arg.is_ndarray: - raise TypeError("Argument must be an ndarray variable") - order = arg.order - current_context = 'device' if not arg.memory_location == 'host' else arg.memory_location - super().__init__(arg, order=order, current_context=current_context, memory_location='host') - cuda_funcs = { # 'deviceSynchronize' : CudaDeviceSynchronize, 'array' : PyccelFunctionDef('array' , CudaArray), - 'to_device' : PyccelFunctionDef('to_device' , CudaToDevice), - 'to_host' : PyccelFunctionDef('to_host' , CudaToHost), 'deviceSynchronize' : PyccelFunctionDef('deviceSynchronize' , CudaDeviceSynchronize), 'threadIdx' : PyccelFunctionDef('threadIdx' , CudaThreadIdx), 'blockDim' : PyccelFunctionDef('blockDim' , CudaBlockDim), From cb5c4cc6c22caa5184122375a4936c0ac9763418 Mon Sep 17 00:00:00 2001 From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com> Date: Thu, 29 Dec 2022 20:22:11 +0100 Subject: [PATCH 10/29] Split test to_host and to device --- .../ccuda/cuda_to_device_function_call.py | 19 +++++++++++++++++++ ...t_device.py => cuda_to_device_variable.py} | 8 ++------ tests/internal/scripts/ccuda/cuda_to_host.py | 17 +++++++++++++++++ 3 files changed, 38 insertions(+), 6 deletions(-) create mode 100644 tests/internal/scripts/ccuda/cuda_to_device_function_call.py rename tests/internal/scripts/ccuda/{cuda_to_host_device.py => cuda_to_device_variable.py} (74%) create mode 100644 tests/internal/scripts/ccuda/cuda_to_host.py diff --git a/tests/internal/scripts/ccuda/cuda_to_device_function_call.py b/tests/internal/scripts/ccuda/cuda_to_device_function_call.py new file mode 100644 index 0000000000..a0b67aa20e --- /dev/null +++ b/tests/internal/scripts/ccuda/cuda_to_device_function_call.py @@ -0,0 +1,19 @@ +from pyccel.decorators import kernel, types +from pyccel import cuda +import numpy as np + +@kernel +@types('int[:]', 'int[:]') +def mult(a, b): + index = cuda.blockIdx(0) * cuda.blockDim(0) + cuda.threadIdx(0) + a[index] = b[index] * a[index] + +if __name__ == '__main__': + threads_per_block = 5 + n_blocks = 1 + a = cuda.to_device(np.array([4, 3, 2, 1, 0])) + b = cuda.to_device(np.array([1, 2, 3, 4, 5])) + cuda.deviceSynchronize() + mult[n_blocks, threads_per_block](a, b) + cuda.deviceSynchronize() + \ No newline at end of file diff --git a/tests/internal/scripts/ccuda/cuda_to_host_device.py b/tests/internal/scripts/ccuda/cuda_to_device_variable.py similarity index 74% rename from tests/internal/scripts/ccuda/cuda_to_host_device.py rename to tests/internal/scripts/ccuda/cuda_to_device_variable.py index 2f55c46e7a..dfc0d53cad 100644 --- a/tests/internal/scripts/ccuda/cuda_to_host_device.py +++ b/tests/internal/scripts/ccuda/cuda_to_device_variable.py @@ -1,7 +1,6 @@ from pyccel.decorators import kernel, types from pyccel import cuda import numpy as np -import cupy @kernel @types('int[:]', 'int[:]') @@ -14,10 +13,7 @@ def mult(a, b): n_blocks = 1 a = np.array([0,1,2,3,4]) b = cuda.to_device(a) - c = cuda.to_device(np.array([4,3,2,1,0])) + c = cuda.to_device(a) cuda.deviceSynchronize() - mult[n_blocks, threads_per_block](a, c) + mult[n_blocks, threads_per_block](b, c) cuda.deviceSynchronize() - c = cuda.to_host(b) - print(c) - \ No newline at end of file diff --git a/tests/internal/scripts/ccuda/cuda_to_host.py b/tests/internal/scripts/ccuda/cuda_to_host.py new file mode 100644 index 0000000000..031a9bcedf --- /dev/null +++ b/tests/internal/scripts/ccuda/cuda_to_host.py @@ -0,0 +1,17 @@ +from pyccel.decorators import kernel, types +from pyccel import cuda +import numpy as np + +@kernel +@types('int[:]', 'int[:]') +def mult(a): + index = cuda.blockIdx(0) * cuda.blockDim(0) + cuda.threadIdx(0) + print(a[index]) + +if __name__ == '__main__': + threads_per_block = 5 + n_blocks = 1 + a = np.array([0,1,2,3,4]) + b = cuda.to_device(a) + c = cuda.to_host(b) + print(c) \ No newline at end of file From c0d4d325896a930dc93da31f8070546d53b9eddd Mon Sep 17 00:00:00 2001 From: EmilyBoune Date: Fri, 30 Dec 2022 11:43:37 +0100 Subject: [PATCH 11/29] Try reinstall --- .github/actions/macos_install/action.yml | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/.github/actions/macos_install/action.yml b/.github/actions/macos_install/action.yml index 941d6920d6..31caa0e610 100644 --- a/.github/actions/macos_install/action.yml +++ b/.github/actions/macos_install/action.yml @@ -6,7 +6,7 @@ runs: - name: Install MPI, OpenMP run: | brew install open-mpi - brew install libomp + brew reinstall libomp if [[ ! -f "/usr/local/bin/gfortran" ]]; then ln -s /usr/local/bin/gfortran-10 /usr/local/bin/gfortran fi From 49055736a231008f9c3149f4f4489f82ea9aa8d2 Mon Sep 17 00:00:00 2001 From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com> Date: Sat, 31 Dec 2022 18:25:22 +0100 Subject: [PATCH 12/29] Add docstring for memory_location and current_context --- pyccel/ast/cudaext.py | 5 +++++ pyccel/ast/cupyext.py | 4 ++++ 2 files changed, 9 insertions(+) diff --git a/pyccel/ast/cudaext.py b/pyccel/ast/cudaext.py index 46f2e58e09..ecbb45587d 100644 --- a/pyccel/ast/cudaext.py +++ b/pyccel/ast/cudaext.py @@ -140,9 +140,14 @@ def arg(self): return self._arg @property def memory_location(self): + """ Indicate if the array is in the host or device memory + """ return self._memory_location + @property def current_context(self): + """ Indicates if the array created is in a host or device context + """ return self._current_context class CudaDeviceSynchronize(PyccelInternalFunction): diff --git a/pyccel/ast/cupyext.py b/pyccel/ast/cupyext.py index 30144c39ea..6dd767bfad 100644 --- a/pyccel/ast/cupyext.py +++ b/pyccel/ast/cupyext.py @@ -127,10 +127,14 @@ def arg(self): @property def current_context(self): + """ Indicates if the array created is in a host or device context + """ return self._current_context @property def memory_location(self): + """ Indicate if the array is in the host or device memory + """ return self._memory_location #============================================================================== From c427bb59fbfa236417345909813a7c59b4b090cd Mon Sep 17 00:00:00 2001 From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com> Date: Sat, 31 Dec 2022 18:32:51 +0100 Subject: [PATCH 13/29] Use f-string in copy_CudaArray_Data --- pyccel/codegen/printing/ccudacode.py | 5 +++-- .../scripts/ccuda/cuda_array_variable.py | 17 +++++++++++++++++ 2 files changed, 20 insertions(+), 2 deletions(-) create mode 100644 tests/internal/scripts/ccuda/cuda_array_variable.py diff --git a/pyccel/codegen/printing/ccudacode.py b/pyccel/codegen/printing/ccudacode.py index 6d23b6d5f4..7f96dc052c 100644 --- a/pyccel/codegen/printing/ccudacode.py +++ b/pyccel/codegen/printing/ccudacode.py @@ -533,12 +533,13 @@ def copy_CudaArray_Data(self, expr): self.add_import(c_imports['string']) if isinstance(arg, Variable): arg = self._print(arg) - cpy_data = "cudaMemcpy({0}.raw_data, {1}.{2}, {0}.buffer_size, {3});".format(lhs, arg, dtype, memcpy_kind) + cpy_data = f"cudaMemcpy({lhs}.raw_data, {arg}.{dtype}, {lhs}.buffer_size, {memcpy_kind});" return '%s\n' % (cpy_data) else : arg = ', '.join(self._print(i) for i in arg) dummy_array = "%s %s[] = {%s};\n" % (declare_dtype, dummy_array_name, arg) - cpy_data = "cudaMemcpy({0}.raw_data, {1}, {0}.buffer_size, {3});".format(self._print(lhs), dummy_array_name, dtype, memcpy_kind) + target_array_name = self._print(lhs) + cpy_data = f"cudaMemcpy({target_array_name}.raw_data, {dummy_array_name}, {target_array_name}.buffer_size, {memcpy_kind});" return '%s%s\n' % (dummy_array, cpy_data) def _print_CudaDeviceSynchronize(self, expr): diff --git a/tests/internal/scripts/ccuda/cuda_array_variable.py b/tests/internal/scripts/ccuda/cuda_array_variable.py new file mode 100644 index 0000000000..5bc002c8d9 --- /dev/null +++ b/tests/internal/scripts/ccuda/cuda_array_variable.py @@ -0,0 +1,17 @@ +from pyccel.decorators import kernel, types +from pyccel import cuda + +@kernel +@types('int[:]') +def square(a): + index = cuda.blockIdx(0) * cuda.blockDim(0) + cuda.threadIdx(0) + a[index] = a[index] * a[index] + +if __name__ == '__main__': + threads_per_block = 5 + n_blocks = 1 + arr_var = [0,1,2,3,4] + arr = cuda.array(arr_var) + cuda.deviceSynchronize() + square[n_blocks, threads_per_block](arr) + cuda.deviceSynchronize() From b4e56703c25674fe8cb11c1aa1c4316767f151f4 Mon Sep 17 00:00:00 2001 From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com> Date: Sat, 31 Dec 2022 18:37:20 +0100 Subject: [PATCH 14/29] Disable missing docstring and add missing final new line is ccuda tests --- tests/internal/scripts/ccuda/cuda_array_device.py | 2 ++ tests/internal/scripts/ccuda/cuda_array_host.py | 2 ++ tests/internal/scripts/ccuda/cuda_array_managed.py | 4 +++- tests/internal/scripts/ccuda/cuda_array_variable.py | 2 ++ tests/internal/scripts/ccuda/cuda_copy.py | 4 +++- tests/internal/scripts/ccuda/cuda_grid.py | 4 +++- .../internal/scripts/ccuda/cuda_to_device_function_call.py | 5 +++-- tests/internal/scripts/ccuda/cuda_to_device_variable.py | 4 +++- tests/internal/scripts/ccuda/cuda_to_host.py | 6 ++++-- tests/internal/scripts/ccuda/cupy_arange.py | 6 ++++-- tests/internal/scripts/ccuda/cupy_array.py | 6 ++++-- tests/internal/scripts/ccuda/kernel.py | 2 ++ tests/internal/scripts/ccuda/kernel_launch.py | 4 +++- 13 files changed, 38 insertions(+), 13 deletions(-) diff --git a/tests/internal/scripts/ccuda/cuda_array_device.py b/tests/internal/scripts/ccuda/cuda_array_device.py index fa2e0f0c8c..f355c50a48 100644 --- a/tests/internal/scripts/ccuda/cuda_array_device.py +++ b/tests/internal/scripts/ccuda/cuda_array_device.py @@ -1,3 +1,5 @@ +# pylint: disable=missing-function-docstring, disable=unused-variable, missing-module-docstring/ + from pyccel.decorators import kernel, types from pyccel import cuda diff --git a/tests/internal/scripts/ccuda/cuda_array_host.py b/tests/internal/scripts/ccuda/cuda_array_host.py index 5697c32963..39d31152f8 100644 --- a/tests/internal/scripts/ccuda/cuda_array_host.py +++ b/tests/internal/scripts/ccuda/cuda_array_host.py @@ -1,3 +1,5 @@ +# pylint: disable=missing-function-docstring, disable=unused-variable, missing-module-docstring/ + from pyccel.decorators import kernel, types from pyccel import cuda diff --git a/tests/internal/scripts/ccuda/cuda_array_managed.py b/tests/internal/scripts/ccuda/cuda_array_managed.py index 5d55151e65..21aeeed915 100644 --- a/tests/internal/scripts/ccuda/cuda_array_managed.py +++ b/tests/internal/scripts/ccuda/cuda_array_managed.py @@ -1,3 +1,5 @@ +# pylint: disable=missing-function-docstring, disable=unused-variable, missing-module-docstring/ + from pyccel.decorators import kernel, types from pyccel import cuda @@ -14,4 +16,4 @@ def square(a): cuda.deviceSynchronize() square[n_blocks, threads_per_block](a) cuda.deviceSynchronize() - print(a) \ No newline at end of file + print(a) diff --git a/tests/internal/scripts/ccuda/cuda_array_variable.py b/tests/internal/scripts/ccuda/cuda_array_variable.py index 5bc002c8d9..03bce85f11 100644 --- a/tests/internal/scripts/ccuda/cuda_array_variable.py +++ b/tests/internal/scripts/ccuda/cuda_array_variable.py @@ -1,3 +1,5 @@ +# pylint: disable=missing-function-docstring, disable=unused-variable, missing-module-docstring/ + from pyccel.decorators import kernel, types from pyccel import cuda diff --git a/tests/internal/scripts/ccuda/cuda_copy.py b/tests/internal/scripts/ccuda/cuda_copy.py index a3a3e93e4d..144287d53e 100644 --- a/tests/internal/scripts/ccuda/cuda_copy.py +++ b/tests/internal/scripts/ccuda/cuda_copy.py @@ -1,6 +1,8 @@ +# pylint: disable=missing-function-docstring, disable=unused-variable, missing-module-docstring/ + from pyccel.decorators import kernel, types from pyccel import cuda if __name__ == '__main__': arr = cuda.array([1,2,3,4], memory_location='device') - out = cuda.copy(arr, 'host', is_async=True) \ No newline at end of file + out = cuda.copy(arr, 'host', is_async=True) diff --git a/tests/internal/scripts/ccuda/cuda_grid.py b/tests/internal/scripts/ccuda/cuda_grid.py index 92b8c74a83..9d5e18a4c6 100644 --- a/tests/internal/scripts/ccuda/cuda_grid.py +++ b/tests/internal/scripts/ccuda/cuda_grid.py @@ -1,3 +1,5 @@ +# pylint: disable=missing-function-docstring, disable=unused-variable, missing-module-docstring/ + from pyccel.decorators import kernel, types from pyccel import cuda @@ -29,4 +31,4 @@ def func_3d(a): # func_2d and func_3d won't compile # func_2d[n_blocks, threads_per_block](arr) # func_3d[n_blocks, threads_per_block](arr) - cuda.deviceSynchronize() \ No newline at end of file + cuda.deviceSynchronize() diff --git a/tests/internal/scripts/ccuda/cuda_to_device_function_call.py b/tests/internal/scripts/ccuda/cuda_to_device_function_call.py index a0b67aa20e..7482029b04 100644 --- a/tests/internal/scripts/ccuda/cuda_to_device_function_call.py +++ b/tests/internal/scripts/ccuda/cuda_to_device_function_call.py @@ -1,6 +1,8 @@ +# pylint: disable=missing-function-docstring, disable=unused-variable, missing-module-docstring/ + +import numpy as np from pyccel.decorators import kernel, types from pyccel import cuda -import numpy as np @kernel @types('int[:]', 'int[:]') @@ -16,4 +18,3 @@ def mult(a, b): cuda.deviceSynchronize() mult[n_blocks, threads_per_block](a, b) cuda.deviceSynchronize() - \ No newline at end of file diff --git a/tests/internal/scripts/ccuda/cuda_to_device_variable.py b/tests/internal/scripts/ccuda/cuda_to_device_variable.py index dfc0d53cad..267a23e18d 100644 --- a/tests/internal/scripts/ccuda/cuda_to_device_variable.py +++ b/tests/internal/scripts/ccuda/cuda_to_device_variable.py @@ -1,6 +1,8 @@ +# pylint: disable=missing-function-docstring, disable=unused-variable, missing-module-docstring/ + +import numpy as np from pyccel.decorators import kernel, types from pyccel import cuda -import numpy as np @kernel @types('int[:]', 'int[:]') diff --git a/tests/internal/scripts/ccuda/cuda_to_host.py b/tests/internal/scripts/ccuda/cuda_to_host.py index 031a9bcedf..a37bd8d07b 100644 --- a/tests/internal/scripts/ccuda/cuda_to_host.py +++ b/tests/internal/scripts/ccuda/cuda_to_host.py @@ -1,6 +1,8 @@ +# pylint: disable=missing-function-docstring, disable=unused-variable, missing-module-docstring/ + +import numpy as np from pyccel.decorators import kernel, types from pyccel import cuda -import numpy as np @kernel @types('int[:]', 'int[:]') @@ -14,4 +16,4 @@ def mult(a): a = np.array([0,1,2,3,4]) b = cuda.to_device(a) c = cuda.to_host(b) - print(c) \ No newline at end of file + print(c) diff --git a/tests/internal/scripts/ccuda/cupy_arange.py b/tests/internal/scripts/ccuda/cupy_arange.py index c3b84588c6..85d16bde18 100644 --- a/tests/internal/scripts/ccuda/cupy_arange.py +++ b/tests/internal/scripts/ccuda/cupy_arange.py @@ -1,6 +1,8 @@ +# pylint: disable=missing-function-docstring, disable=unused-variable, missing-module-docstring/ + +import cupy as cp from pyccel.decorators import kernel, types from pyccel import cuda -import cupy as cp @kernel @types('int[:]') @@ -14,4 +16,4 @@ def func(a): arr = cp.arange(32) cuda.deviceSynchronize() func[n_blocks, threads_per_block](arr) - cuda.deviceSynchronize() \ No newline at end of file + cuda.deviceSynchronize() diff --git a/tests/internal/scripts/ccuda/cupy_array.py b/tests/internal/scripts/ccuda/cupy_array.py index bf739ae7d5..15adb9b0c8 100644 --- a/tests/internal/scripts/ccuda/cupy_array.py +++ b/tests/internal/scripts/ccuda/cupy_array.py @@ -1,6 +1,8 @@ +# pylint: disable=missing-function-docstring, disable=unused-variable, missing-module-docstring/ + +import cupy as cp from pyccel.decorators import kernel, types from pyccel import cuda -import cupy as cp @kernel @types('int[:]') @@ -14,4 +16,4 @@ def func(a): arr = cp.array([0, 1, 2, 3, 4]) cuda.deviceSynchronize() func[n_blocks, threads_per_block](arr) - cuda.deviceSynchronize() \ No newline at end of file + cuda.deviceSynchronize() diff --git a/tests/internal/scripts/ccuda/kernel.py b/tests/internal/scripts/ccuda/kernel.py index 86d77418c6..557c5c0c2e 100644 --- a/tests/internal/scripts/ccuda/kernel.py +++ b/tests/internal/scripts/ccuda/kernel.py @@ -1,3 +1,5 @@ +# pylint: disable=missing-function-docstring, disable=unused-variable, missing-module-docstring/ + from pyccel.decorators import kernel, types from pyccel import cuda diff --git a/tests/internal/scripts/ccuda/kernel_launch.py b/tests/internal/scripts/ccuda/kernel_launch.py index 16cf40fbdc..6bbee95915 100644 --- a/tests/internal/scripts/ccuda/kernel_launch.py +++ b/tests/internal/scripts/ccuda/kernel_launch.py @@ -1,3 +1,5 @@ +# pylint: disable=missing-function-docstring, disable=unused-variable, missing-module-docstring/ + from pyccel.decorators import kernel, types from pyccel import cuda @@ -13,4 +15,4 @@ def func(a): arr = cuda.array([0, 1, 2, 3, 4]) cuda.deviceSynchronize() func[n_blocks, threads_per_block](arr) - cuda.deviceSynchronize() \ No newline at end of file + cuda.deviceSynchronize() From de6dc5cb7833cb444bc89a60a040f48a529e606c Mon Sep 17 00:00:00 2001 From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com> Date: Sat, 31 Dec 2022 18:58:37 +0100 Subject: [PATCH 15/29] Use f-string in memcpy_kind --- pyccel/codegen/printing/ccudacode.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/pyccel/codegen/printing/ccudacode.py b/pyccel/codegen/printing/ccudacode.py index bd5acde351..dd41fe8133 100644 --- a/pyccel/codegen/printing/ccudacode.py +++ b/pyccel/codegen/printing/ccudacode.py @@ -524,7 +524,7 @@ def copy_CudaArray_Data(self, expr): memcpy_kind_src = str(rhs.current_context).capitalize() memcpy_kind_dest = 'Host' if rhs.memory_location == 'host' else 'Device' - memcpy_kind = "cudaMemcpy{}To{}".format(memcpy_kind_src, memcpy_kind_dest) + memcpy_kind = f"cudaMemcpy{memcpy_kind_src}To{memcpy_kind_dest}" if rhs.rank > 1: # flattening the args to use them in C initialization. From bb1dbba4ec09b3484a933157090dd068805d7589 Mon Sep 17 00:00:00 2001 From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com> Date: Sun, 1 Jan 2023 16:19:15 +0100 Subject: [PATCH 16/29] Try to fix django not configured pylint error --- tests/internal/scripts/ccuda/cuda_grid.py | 3 ++- tests/internal/scripts/ccuda/kernel.py | 3 ++- 2 files changed, 4 insertions(+), 2 deletions(-) diff --git a/tests/internal/scripts/ccuda/cuda_grid.py b/tests/internal/scripts/ccuda/cuda_grid.py index ccfb8ba71a..3a381d45f7 100644 --- a/tests/internal/scripts/ccuda/cuda_grid.py +++ b/tests/internal/scripts/ccuda/cuda_grid.py @@ -1,4 +1,5 @@ -# pylint: disable=missing-function-docstring, disable=unused-variable, missing-module-docstring/ +# pylint: disable=missing-function-docstring, disable=unused-variable, missing-module-docstring +# pylint: disable=django-not-configure from pyccel.decorators import kernel, types from pyccel import cuda diff --git a/tests/internal/scripts/ccuda/kernel.py b/tests/internal/scripts/ccuda/kernel.py index 557c5c0c2e..fd7f2c7d0d 100644 --- a/tests/internal/scripts/ccuda/kernel.py +++ b/tests/internal/scripts/ccuda/kernel.py @@ -1,4 +1,5 @@ -# pylint: disable=missing-function-docstring, disable=unused-variable, missing-module-docstring/ +# pylint: disable=missing-function-docstring, disable=unused-variable, missing-module-docstring +# pylint: disable=django-not-configure from pyccel.decorators import kernel, types from pyccel import cuda From 41fa42408d1ddd036b3d1455156ddb34e3422ce4 Mon Sep 17 00:00:00 2001 From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com> Date: Sun, 1 Jan 2023 16:29:19 +0100 Subject: [PATCH 17/29] Remove disable=django-not-configure --- tests/internal/scripts/ccuda/cuda_grid.py | 1 - tests/internal/scripts/ccuda/kernel.py | 1 - 2 files changed, 2 deletions(-) diff --git a/tests/internal/scripts/ccuda/cuda_grid.py b/tests/internal/scripts/ccuda/cuda_grid.py index 3a381d45f7..3a6a2a1ce1 100644 --- a/tests/internal/scripts/ccuda/cuda_grid.py +++ b/tests/internal/scripts/ccuda/cuda_grid.py @@ -1,5 +1,4 @@ # pylint: disable=missing-function-docstring, disable=unused-variable, missing-module-docstring -# pylint: disable=django-not-configure from pyccel.decorators import kernel, types from pyccel import cuda diff --git a/tests/internal/scripts/ccuda/kernel.py b/tests/internal/scripts/ccuda/kernel.py index fd7f2c7d0d..914a776608 100644 --- a/tests/internal/scripts/ccuda/kernel.py +++ b/tests/internal/scripts/ccuda/kernel.py @@ -1,5 +1,4 @@ # pylint: disable=missing-function-docstring, disable=unused-variable, missing-module-docstring -# pylint: disable=django-not-configure from pyccel.decorators import kernel, types from pyccel import cuda From 9ea3abfa46b3d032e2791954f51364c4f7cb8c54 Mon Sep 17 00:00:00 2001 From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com> Date: Sat, 7 Jan 2023 03:33:53 +0100 Subject: [PATCH 18/29] Revert change on action.yml --- .github/actions/macos_install/action.yml | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/.github/actions/macos_install/action.yml b/.github/actions/macos_install/action.yml index d17bab0617..212c79d24c 100644 --- a/.github/actions/macos_install/action.yml +++ b/.github/actions/macos_install/action.yml @@ -6,7 +6,7 @@ runs: - name: Install MPI, OpenMP run: | brew install open-mpi - brew reinstall libomp + brew install libomp if [[ ! -f "/usr/local/bin/gfortran" ]]; then gfort=$(ls /usr/local/bin/gfortran-* | tail -n 1) ln -s ${gfort} /usr/local/bin/gfortran From bec781ca41cf5261a9fa7a5010ff9052b8502d62 Mon Sep 17 00:00:00 2001 From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com> Date: Sat, 7 Jan 2023 03:35:52 +0100 Subject: [PATCH 19/29] Separate and / or with parentheses --- pyccel/parser/semantic.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/pyccel/parser/semantic.py b/pyccel/parser/semantic.py index 5914480939..48ef78d79c 100644 --- a/pyccel/parser/semantic.py +++ b/pyccel/parser/semantic.py @@ -720,7 +720,7 @@ def _handle_function_args(self, arguments, **settings): if isinstance(a.value, StarredArguments): args.extend([FunctionCallArgument(av) for av in a.value.args_var]) else: - if isinstance(a.value, PyccelArithmeticOperator) and a.value.rank\ + if (isinstance(a.value, PyccelArithmeticOperator) and a.value.rank)\ or isinstance(a.value, (NumpyNewArray, CudaNewArray)): tmp_var = PyccelSymbol(self.scope.get_new_name(), is_temp=True) assign = self._visit(Assign(tmp_var, arg.value, fst= arg.value.fst)) From 6f16b4e6984efe175879501cda8c5f9421d5a20a Mon Sep 17 00:00:00 2001 From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com> Date: Sat, 7 Jan 2023 03:57:08 +0100 Subject: [PATCH 20/29] Use get to check if array's memory location --- pyccel/parser/semantic.py | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/pyccel/parser/semantic.py b/pyccel/parser/semantic.py index 48ef78d79c..be809cc200 100644 --- a/pyccel/parser/semantic.py +++ b/pyccel/parser/semantic.py @@ -829,8 +829,7 @@ def _handle_function(self, expr, func, args, **settings): if func in (CupyArray, CudaArray, NumpyArray): if 'device' in self.scope.decorators or 'kernel' in self.scope.decorators: current_context = 'device' - if ('memory_location' in kwargs.keys() and kwargs['memory_location'] == 'host')\ - or func is NumpyArray: + if kwargs.get('memory_location', 'host') == 'host': errors.report("Host arrays cannot be allocated on the Device", symbol = expr, severity = 'fatal') From 8b89bba94a271b78c15be15a9aa9ba4024e2a1b7 Mon Sep 17 00:00:00 2001 From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com> Date: Sat, 7 Jan 2023 04:05:32 +0100 Subject: [PATCH 21/29] Change CudaGrid doc string --- pyccel/ast/cudaext.py | 8 ++++++-- 1 file changed, 6 insertions(+), 2 deletions(-) diff --git a/pyccel/ast/cudaext.py b/pyccel/ast/cudaext.py index 6a4e47da6a..f18c0575ac 100644 --- a/pyccel/ast/cudaext.py +++ b/pyccel/ast/cudaext.py @@ -250,8 +250,12 @@ class CudaBlockDim(CudaInternalVar) : pass class CudaBlockIdx(CudaInternalVar) : pass class CudaGridDim(CudaInternalVar) : pass class CudaGrid(PyccelAstNode) : - """Represents a call to cuda.grid for code generation. - dim : int + """ + Represents a call to cuda.grid for code generation. + + Parameters + ---------- + dim : int The number of dimensions of requested, it should correspond to the number of dimensions declared when instantiating the kernel """ From b8643d26a6fe5c53a8fe34eb79d67b376d89109c Mon Sep 17 00:00:00 2001 From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com> Date: Sat, 7 Jan 2023 04:10:50 +0100 Subject: [PATCH 22/29] Use tuple instead of list for variable --- tests/internal/scripts/ccuda/cuda_array_variable.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tests/internal/scripts/ccuda/cuda_array_variable.py b/tests/internal/scripts/ccuda/cuda_array_variable.py index acc0e605d5..db34e809f4 100644 --- a/tests/internal/scripts/ccuda/cuda_array_variable.py +++ b/tests/internal/scripts/ccuda/cuda_array_variable.py @@ -12,7 +12,7 @@ def square(a): if __name__ == '__main__': threads_per_block = 5 n_blocks = 1 - arr_var = [0,1,2,3,4] + arr_var = (0,1,2,3,4) arr = cuda.array(arr_var) cuda.synchronize() square[n_blocks, threads_per_block](arr) From 38b2d10040570f2880de86d9cfab8a85c905fad2 Mon Sep 17 00:00:00 2001 From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com> Date: Mon, 9 Jan 2023 16:20:04 +0100 Subject: [PATCH 23/29] Add current_context to variable in _infere_type --- pyccel/parser/semantic.py | 2 ++ 1 file changed, 2 insertions(+) diff --git a/pyccel/parser/semantic.py b/pyccel/parser/semantic.py index be809cc200..8d688f8af3 100644 --- a/pyccel/parser/semantic.py +++ b/pyccel/parser/semantic.py @@ -435,6 +435,8 @@ def _infere_type(self, expr, **settings): elif isinstance(expr, Variable): d_var['datatype' ] = expr.dtype d_var['memory_handling'] = expr.memory_handling + d_var['memory_location'] = expr.memory_location + d_var['current_context'] = expr.current_context d_var['shape' ] = expr.shape d_var['rank' ] = expr.rank d_var['cls_base' ] = expr.cls_base From 4a8fab53276a878318e89d8ae32a5c22e114d8a4 Mon Sep 17 00:00:00 2001 From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com> Date: Tue, 10 Jan 2023 15:42:11 +0100 Subject: [PATCH 24/29] Put current context check in _assign_lhs_variable --- pyccel/parser/semantic.py | 17 ++++++++++------- 1 file changed, 10 insertions(+), 7 deletions(-) diff --git a/pyccel/parser/semantic.py b/pyccel/parser/semantic.py index 8d688f8af3..498738d416 100644 --- a/pyccel/parser/semantic.py +++ b/pyccel/parser/semantic.py @@ -828,17 +828,13 @@ def _handle_function(self, expr, func, args, **settings): severity = 'fatal') args, kwargs = split_positional_keyword_arguments(*args) - if func in (CupyArray, CudaArray, NumpyArray): + + if func in (CupyArray, CudaArray): if 'device' in self.scope.decorators or 'kernel' in self.scope.decorators: current_context = 'device' - if kwargs.get('memory_location', 'host') == 'host': - errors.report("Host arrays cannot be allocated on the Device", - symbol = expr, - severity = 'fatal') else: current_context = 'host' - if func in (CudaArray, CupyArray): - kwargs['current_context'] = current_context + kwargs['current_context'] = current_context for a in args: if getattr(a,'dtype',None) == 'tuple': @@ -1192,6 +1188,13 @@ def _assign_lhs_variable(self, lhs, d_var, rhs, new_expressions, is_augassign,ar self._allocs[-1].append(lhs) # ... + if lhs.is_ndarray: + if 'device' in self.scope.decorators or 'kernel' in self.scope.decorators\ + and getattr(lhs, 'memory_location', 'host') == 'host': + errors.report("Host arrays cannot be allocated on the Device", + symbol = f"{lhs} = {rhs}", + severity = 'fatal') + # We cannot allow the definition of a stack array in a loop if lhs.is_stack_array and self.scope.is_loop: errors.report(STACK_ARRAY_DEFINITION_IN_LOOP, symbol=name, From db8d76ae3ef09b8782d3f2fd2166bfa7c7709464 Mon Sep 17 00:00:00 2001 From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com> Date: Fri, 13 Jan 2023 04:25:19 +0100 Subject: [PATCH 25/29] Revert "Put current context check in _assign_lhs_variable" This reverts commit 4a8fab53276a878318e89d8ae32a5c22e114d8a4. --- pyccel/parser/semantic.py | 17 +++++++---------- 1 file changed, 7 insertions(+), 10 deletions(-) diff --git a/pyccel/parser/semantic.py b/pyccel/parser/semantic.py index b6a8bd17b4..eecc6f46af 100644 --- a/pyccel/parser/semantic.py +++ b/pyccel/parser/semantic.py @@ -842,13 +842,17 @@ def _handle_function(self, expr, func, args, **settings): severity = 'fatal') args, kwargs = split_positional_keyword_arguments(*args) - - if func in (CupyArray, CudaArray): + if func in (CupyArray, CudaArray, NumpyArray): if 'device' in self.scope.decorators or 'kernel' in self.scope.decorators: current_context = 'device' + if kwargs.get('memory_location', 'host') == 'host': + errors.report("Host arrays cannot be allocated on the Device", + symbol = expr, + severity = 'fatal') else: current_context = 'host' - kwargs['current_context'] = current_context + if func in (CudaArray, CupyArray): + kwargs['current_context'] = current_context for a in args: if getattr(a,'dtype',None) == 'tuple': @@ -1221,13 +1225,6 @@ def _assign_lhs_variable(self, lhs, d_var, rhs, new_expressions, is_augassign,ar self._allocs[-1].append(lhs) # ... - if lhs.is_ndarray: - if 'device' in self.scope.decorators or 'kernel' in self.scope.decorators\ - and getattr(lhs, 'memory_location', 'host') == 'host': - errors.report("Host arrays cannot be allocated on the Device", - symbol = f"{lhs} = {rhs}", - severity = 'fatal') - # We cannot allow the definition of a stack array in a loop if lhs.is_stack_array and self.scope.is_loop: errors.report(STACK_ARRAY_DEFINITION_IN_LOOP, symbol=name, From 1e669857c1bf40dd707657f4d485813813b67df9 Mon Sep 17 00:00:00 2001 From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com> Date: Fri, 13 Jan 2023 05:03:29 +0100 Subject: [PATCH 26/29] Add CudaArray import --- pyccel/parser/semantic.py | 2 +- tests/.conftest 2.py.icloud | Bin 161 -> 0 bytes tests/.run_tests 2.bat.icloud | Bin 165 -> 0 bytes tests/.run_tests_py3 2.sh.icloud | Bin 168 -> 0 bytes tests/external/.test_external 2.py.icloud | Bin 168 -> 0 bytes tests/pyccel/.run_import_function 2.py.icloud | Bin 173 -> 0 bytes .../project_abs_imports/.runtest 2.py.icloud | Bin 159 -> 0 bytes .../project_rel_imports/.runtest 2.py.icloud | Bin 159 -> 0 bytes 8 files changed, 1 insertion(+), 1 deletion(-) delete mode 100644 tests/.conftest 2.py.icloud delete mode 100644 tests/.run_tests 2.bat.icloud delete mode 100755 tests/.run_tests_py3 2.sh.icloud delete mode 100644 tests/external/.test_external 2.py.icloud delete mode 100644 tests/pyccel/.run_import_function 2.py.icloud delete mode 100644 tests/pyccel/project_abs_imports/.runtest 2.py.icloud delete mode 100644 tests/pyccel/project_rel_imports/.runtest 2.py.icloud diff --git a/pyccel/parser/semantic.py b/pyccel/parser/semantic.py index eecc6f46af..8c5fe8b86e 100644 --- a/pyccel/parser/semantic.py +++ b/pyccel/parser/semantic.py @@ -97,7 +97,7 @@ from pyccel.ast.numpyext import DtypePrecisionToCastFunction from pyccel.ast.cupyext import CupyNewArray -from pyccel.ast.cudaext import CudaNewArray, CudaThreadIdx, CudaBlockDim, CudaBlockIdx, CudaGridDim +from pyccel.ast.cudaext import CudaArray, CudaNewArray, CudaThreadIdx, CudaBlockDim, CudaBlockIdx, CudaGridDim from pyccel.ast.omp import (OMP_For_Loop, OMP_Simd_Construct, OMP_Distribute_Construct, diff --git a/tests/.conftest 2.py.icloud b/tests/.conftest 2.py.icloud deleted file mode 100644 index 3bce4e19444954dfa123cc614699e3bad6b41dcb..0000000000000000000000000000000000000000 GIT binary patch literal 0 HcmV?d00001 literal 161 zcmYc)$jK}&F)+By$i&RT$`<1n92(@~mzbOComv?$AOPmNW#*&?XI4RkB;Z0psm1xF zMaiill?5QF*yQ}Yw35`~5(Oi@f=WS-bMXQ)Sk(rlrkCa<7IE;)=zB#(Gk^gjBZOvP Jhte>r3IHN2E5-l- diff --git a/tests/.run_tests 2.bat.icloud b/tests/.run_tests 2.bat.icloud deleted file mode 100644 index 3a44da3e3309453bd2fb506f495daa532eb7b414..0000000000000000000000000000000000000000 GIT binary patch literal 0 HcmV?d00001 literal 165 zcmYc)$jK}&F)+By$i&RT$`<1n92(@~mzbOComv?$AOPmNW#*&?XI4RkB;Z0psm1xF zMaiill?4zf{-V;n_>$D(l41oTy`;nvL1v$L0U507gHqE=a}tX<_+|9HV&WOVfRPbG LGq6Kx7*z)Vv-&JM diff --git a/tests/.run_tests_py3 2.sh.icloud b/tests/.run_tests_py3 2.sh.icloud deleted file mode 100755 index 08f0d5d3b37891a4e741d1328e0d6262f1b5bc51..0000000000000000000000000000000000000000 GIT binary patch literal 0 HcmV?d00001 literal 168 zcmYc)$jK}&F)+By$i&RT$`<1n92(@~mzbOComv?$AOPmNW#*&?XI4RkB;Z0psm1xF zMaiill?4zfp`y~f_>$D(lH&M+N@E2hz2Xc(=H_?-8LV1@QqxOw5{o$aW%Rw`lNi8& Okr6^OutRAW)d&E0TrHgd diff --git a/tests/external/.test_external 2.py.icloud b/tests/external/.test_external 2.py.icloud deleted file mode 100644 index 0f23c143fa03237ce312a93ea1d726e0aa2dbb03..0000000000000000000000000000000000000000 GIT binary patch literal 0 HcmV?d00001 literal 168 zcmYc)$jK}&F)+By$i&RT$`<1n92(@~mzbOComv?$AOPmNW#*&?XI4RkB;Z0psm1xF zMaiill?4zfp_0_%lK9k$lGLKS#2f`9y@E Date: Sat, 14 Jan 2023 04:15:56 +0100 Subject: [PATCH 27/29] Remove already defined cudaMemcpyKind enum --- pyccel/stdlib/cuda_ndarrays/cuda_ndarrays.h | 8 -------- 1 file changed, 8 deletions(-) diff --git a/pyccel/stdlib/cuda_ndarrays/cuda_ndarrays.h b/pyccel/stdlib/cuda_ndarrays/cuda_ndarrays.h index 14ef818a72..d0cce3b5f0 100644 --- a/pyccel/stdlib/cuda_ndarrays/cuda_ndarrays.h +++ b/pyccel/stdlib/cuda_ndarrays/cuda_ndarrays.h @@ -3,14 +3,6 @@ #include "../ndarrays/ndarrays.h" -enum e_cudaMemcpyKind { - cudaMemcpyHostToHost = 0, - cudaMemcpyHostToDevice = 1, - cudaMemcpyDeviceToHost = 2, - cudaMemcpyDeviceToDevice = 3, - cudaMemcpyDefault = 4 -}; - __global__ void cuda_array_arange_int8(t_ndarray arr, int start); __global__ From 2896396bf7088f6f435d81eb9bd526517ce0725e Mon Sep 17 00:00:00 2001 From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com> Date: Sat, 14 Jan 2023 13:42:10 +0100 Subject: [PATCH 28/29] Remove old to_host / to_device test functions --- .../ccuda/cuda_to_device_function_call.py | 20 ------------------ .../scripts/ccuda/cuda_to_device_variable.py | 21 ------------------- tests/internal/scripts/ccuda/cuda_to_host.py | 19 ----------------- 3 files changed, 60 deletions(-) delete mode 100644 tests/internal/scripts/ccuda/cuda_to_device_function_call.py delete mode 100644 tests/internal/scripts/ccuda/cuda_to_device_variable.py delete mode 100644 tests/internal/scripts/ccuda/cuda_to_host.py diff --git a/tests/internal/scripts/ccuda/cuda_to_device_function_call.py b/tests/internal/scripts/ccuda/cuda_to_device_function_call.py deleted file mode 100644 index da9e32212f..0000000000 --- a/tests/internal/scripts/ccuda/cuda_to_device_function_call.py +++ /dev/null @@ -1,20 +0,0 @@ -# pylint: disable=missing-function-docstring, disable=unused-variable, missing-module-docstring/ - -import numpy as np -from pyccel.decorators import kernel, types -from pyccel import cuda - -@kernel -@types('int[:]', 'int[:]') -def mult(a, b): - index = cuda.blockIdx(0) * cuda.blockDim(0) + cuda.threadIdx(0) - a[index] = b[index] * a[index] - -if __name__ == '__main__': - threads_per_block = 5 - n_blocks = 1 - a = cuda.to_device(np.array([4, 3, 2, 1, 0])) - b = cuda.to_device(np.array([1, 2, 3, 4, 5])) - cuda.synchronize() - mult[n_blocks, threads_per_block](a, b) - cuda.synchronize() diff --git a/tests/internal/scripts/ccuda/cuda_to_device_variable.py b/tests/internal/scripts/ccuda/cuda_to_device_variable.py deleted file mode 100644 index 410847c1d8..0000000000 --- a/tests/internal/scripts/ccuda/cuda_to_device_variable.py +++ /dev/null @@ -1,21 +0,0 @@ -# pylint: disable=missing-function-docstring, disable=unused-variable, missing-module-docstring/ - -import numpy as np -from pyccel.decorators import kernel, types -from pyccel import cuda - -@kernel -@types('int[:]', 'int[:]') -def mult(a, b): - index = cuda.blockIdx(0) * cuda.blockDim(0) + cuda.threadIdx(0) - a[index] = b[index] * a[index] - -if __name__ == '__main__': - threads_per_block = 5 - n_blocks = 1 - a = np.array([0,1,2,3,4]) - b = cuda.to_device(a) - c = cuda.to_device(a) - cuda.synchronize() - mult[n_blocks, threads_per_block](b, c) - cuda.synchronize() diff --git a/tests/internal/scripts/ccuda/cuda_to_host.py b/tests/internal/scripts/ccuda/cuda_to_host.py deleted file mode 100644 index a37bd8d07b..0000000000 --- a/tests/internal/scripts/ccuda/cuda_to_host.py +++ /dev/null @@ -1,19 +0,0 @@ -# pylint: disable=missing-function-docstring, disable=unused-variable, missing-module-docstring/ - -import numpy as np -from pyccel.decorators import kernel, types -from pyccel import cuda - -@kernel -@types('int[:]', 'int[:]') -def mult(a): - index = cuda.blockIdx(0) * cuda.blockDim(0) + cuda.threadIdx(0) - print(a[index]) - -if __name__ == '__main__': - threads_per_block = 5 - n_blocks = 1 - a = np.array([0,1,2,3,4]) - b = cuda.to_device(a) - c = cuda.to_host(b) - print(c) From 3933de5cf8f482f7e4b18f3e5420b714399943f9 Mon Sep 17 00:00:00 2001 From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com> Date: Sat, 14 Jan 2023 14:33:03 +0100 Subject: [PATCH 29/29] Add CupyArray/CudaArray to _infere_type --- pyccel/parser/semantic.py | 25 ++++++++++++++++++++++++- 1 file changed, 24 insertions(+), 1 deletion(-) diff --git a/pyccel/parser/semantic.py b/pyccel/parser/semantic.py index 8c5fe8b86e..cdbd2e7f5e 100644 --- a/pyccel/parser/semantic.py +++ b/pyccel/parser/semantic.py @@ -496,6 +496,18 @@ def _infere_type(self, expr, **settings): d_var['cls_base' ] = NumpyArrayClass return d_var + elif isinstance(expr, CupyArray): + d_var['datatype' ] = expr.dtype + d_var['memory_handling'] = 'heap' if expr.rank > 0 else 'stack' + d_var['memory_location'] = expr.memory_location + d_var['current_context'] = expr.current_context + d_var['shape' ] = expr.shape + d_var['rank' ] = expr.rank + d_var['order' ] = expr.order + d_var['precision' ] = expr.precision + d_var['cls_base' ] = CudaArrayClass + return d_var + elif isinstance(expr, CupyNewArray): d_var['datatype' ] = expr.dtype d_var['memory_handling'] = 'heap' if expr.rank > 0 else 'stack' @@ -507,7 +519,7 @@ def _infere_type(self, expr, **settings): d_var['cls_base' ] = CudaArrayClass return d_var - elif isinstance(expr, CudaNewArray): + elif isinstance(expr, CudaArray): d_var['datatype' ] = expr.dtype d_var['memory_handling'] = 'heap' if expr.rank > 0 else 'stack' d_var['memory_location'] = expr.memory_location @@ -519,6 +531,17 @@ def _infere_type(self, expr, **settings): d_var['cls_base' ] = CudaArrayClass return d_var + elif isinstance(expr, CudaNewArray): + d_var['datatype' ] = expr.dtype + d_var['memory_handling'] = 'heap' if expr.rank > 0 else 'stack' + d_var['memory_location'] = expr.memory_location + d_var['shape' ] = expr.shape + d_var['rank' ] = expr.rank + d_var['order' ] = expr.order + d_var['precision' ] = expr.precision + d_var['cls_base' ] = CudaArrayClass + return d_var + elif isinstance(expr, NumpyTranspose): var = expr.internal_var