From ba7187423aabcdd824c62381e811ddc570d35f55 Mon Sep 17 00:00:00 2001
From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com>
Date: Wed, 26 Oct 2022 16:54:20 +0100
Subject: [PATCH 01/32] Add cupy rabel

---
 pyccel/ast/cupyext.py                | 28 ++++++++++++++++++++++++++++
 pyccel/codegen/printing/ccudacode.py |  8 ++++----
 2 files changed, 32 insertions(+), 4 deletions(-)

diff --git a/pyccel/ast/cupyext.py b/pyccel/ast/cupyext.py
index 08650db3eb..8e579aa6d7 100644
--- a/pyccel/ast/cupyext.py
+++ b/pyccel/ast/cupyext.py
@@ -55,6 +55,7 @@
     'Shape',
     'CupyZeros',
     'CupyZerosLike',
+    'CupyRavel'
 )
 
 #==============================================================================
@@ -436,6 +437,32 @@ def __str__(self):
 
 #==============================================================================
 
+class CupyRavel(CupyArray):
+
+    name   = 'ravel'
+    __slots__ = ('_arg','_dtype','_precision','_shape','_rank','_order', '_memory_location')
+    _attribute_nodes = ('_arg',)
+    def __new__(cls, arg):
+        if not isinstance(arg, (list, tuple, PyccelAstNode)):
+            raise TypeError('Unknown type of  %s.' % type(arg))
+        elif arg.rank < 2:
+            return arg
+        else:
+            return super().__new__(cls)
+
+    def __init__(self, arg, memory_location='managed'):
+        self._arg = arg
+        super().__init__(arg)
+        self._shape = [LiteralInteger(reduce((lambda x, y: x.python_value * y.python_value), self.shape))]
+        self._rank = len(self._shape)
+        self._order = None
+
+
+    @property
+    def arg(self):
+        return self._arg
+
+#==============================================================================
 cupy_funcs = {
     # ... array creation routines
     'full'      : PyccelFunctionDef('full'      , CupyFull),
@@ -448,6 +475,7 @@ def __str__(self):
     'ones_like' : PyccelFunctionDef('ones_like' , CupyOnesLike),
     'array'     : PyccelFunctionDef('array'     , CupyArray),
     'arange'    : PyccelFunctionDef('arange'    , CupyArange),
+    'ravel'     : PyccelFunctionDef('ravel'     , CupyRavel),
     # ...
     'shape'     : PyccelFunctionDef('shape'     , Shape),
     'size'      : PyccelFunctionDef('size'      , CupyArraySize),
diff --git a/pyccel/codegen/printing/ccudacode.py b/pyccel/codegen/printing/ccudacode.py
index fcf7d62eb2..4865714dd4 100644
--- a/pyccel/codegen/printing/ccudacode.py
+++ b/pyccel/codegen/printing/ccudacode.py
@@ -40,7 +40,7 @@
 from pyccel.ast.numpyext import NumpyFull, NumpyArray, NumpyArange
 from pyccel.ast.numpyext import NumpyReal, NumpyImag, NumpyFloat
 
-from pyccel.ast.cupyext import CupyFull, CupyArray, CupyArange
+from pyccel.ast.cupyext import CupyFull, CupyArray, CupyArange, CupyRavel
 
 from pyccel.ast.cudaext import cuda_Internal_Var, CudaArray
 
@@ -521,9 +521,6 @@ def copy_CudaArray_Data(self, expr):
         declare_dtype = self.find_in_dtype_registry(self._print(rhs.dtype), rhs.precision)
         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.rank > 1:
-            # flattening the args to use them in C initialization.
-            arg = self._flatten_list(arg)
 
         self.add_import(c_imports['string'])
         if isinstance(arg, Variable):
@@ -531,6 +528,9 @@ def copy_CudaArray_Data(self, expr):
             cpy_data = "cudaMemcpy({0}.raw_data, {1}.{2}, {0}.buffer_size, cudaMemcpyHostToDevice);".format(lhs, arg, dtype)
             return '%s\n' % (cpy_data)
         else :
+            if arg.rank > 1:
+            # flattening the args to use them in C initialization.
+                arg = self._flatten_list(arg)
             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)

From 8207c2db4c33fafcbc3bdd7ed049521590b0c00a Mon Sep 17 00:00:00 2001
From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com>
Date: Wed, 26 Oct 2022 17:03:11 +0100
Subject: [PATCH 02/32] Add docstring for ravel

---
 pyccel/ast/cupyext.py | 11 ++++++++++-
 1 file changed, 10 insertions(+), 1 deletion(-)

diff --git a/pyccel/ast/cupyext.py b/pyccel/ast/cupyext.py
index 8e579aa6d7..459cbee6ef 100644
--- a/pyccel/ast/cupyext.py
+++ b/pyccel/ast/cupyext.py
@@ -438,7 +438,17 @@ def __str__(self):
 #==============================================================================
 
 class CupyRavel(CupyArray):
+    """
+    Class representing a call to the cupy ravel function which
+    returns flattened version of the passed array
 
+    Parameters
+    ==========
+    arg : PyccelAstNode
+            A PyccelAstNode of unknown shape
+    memory_location : string
+            The location where the new array memory should be allocated
+    """
     name   = 'ravel'
     __slots__ = ('_arg','_dtype','_precision','_shape','_rank','_order', '_memory_location')
     _attribute_nodes = ('_arg',)
@@ -457,7 +467,6 @@ def __init__(self, arg, memory_location='managed'):
         self._rank = len(self._shape)
         self._order = None
 
-
     @property
     def arg(self):
         return self._arg

From d9490863c7ae301aa864a2b6ed1334de7aa4dfe3 Mon Sep 17 00:00:00 2001
From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com>
Date: Wed, 26 Oct 2022 17:41:56 +0100
Subject: [PATCH 03/32] fix cuda copy array printing

---
 pyccel/ast/cupyext.py                | 9 +++------
 pyccel/codegen/printing/ccudacode.py | 2 +-
 2 files changed, 4 insertions(+), 7 deletions(-)

diff --git a/pyccel/ast/cupyext.py b/pyccel/ast/cupyext.py
index 459cbee6ef..2c1c3f348c 100644
--- a/pyccel/ast/cupyext.py
+++ b/pyccel/ast/cupyext.py
@@ -455,15 +455,12 @@ class CupyRavel(CupyArray):
     def __new__(cls, arg):
         if not isinstance(arg, (list, tuple, PyccelAstNode)):
             raise TypeError('Unknown type of  %s.' % type(arg))
-        elif arg.rank < 2:
-            return arg
-        else:
-            return super().__new__(cls)
+        return super().__new__(cls)
 
     def __init__(self, arg, memory_location='managed'):
-        self._arg = arg
         super().__init__(arg)
-        self._shape = [LiteralInteger(reduce((lambda x, y: x.python_value * y.python_value), self.shape))]
+        shape = reduce((lambda x, y: x.python_value * y.python_value), self.shape)
+        self._shape = [shape if isinstance(shape, LiteralInteger) else LiteralInteger(shape)]
         self._rank = len(self._shape)
         self._order = None
 
diff --git a/pyccel/codegen/printing/ccudacode.py b/pyccel/codegen/printing/ccudacode.py
index 4865714dd4..0c087fc690 100644
--- a/pyccel/codegen/printing/ccudacode.py
+++ b/pyccel/codegen/printing/ccudacode.py
@@ -525,7 +525,7 @@ 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}.raw_data, {0}.buffer_size, cudaMemcpyHostToDevice);".format(lhs, arg)
             return '%s\n' % (cpy_data)
         else :
             if arg.rank > 1:

From 4d56a590d34ea4bdb45834461c98c1d2fa8853c1 Mon Sep 17 00:00:00 2001
From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com>
Date: Wed, 26 Oct 2022 17:52:40 +0100
Subject: [PATCH 04/32] Fix errors in Kernel Semantic checks

---
 pyccel/parser/semantic.py | 6 +++---
 1 file changed, 3 insertions(+), 3 deletions(-)

diff --git a/pyccel/parser/semantic.py b/pyccel/parser/semantic.py
index 023effaff8..e46a957eec 100644
--- a/pyccel/parser/semantic.py
+++ b/pyccel/parser/semantic.py
@@ -812,7 +812,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 +902,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 8719440854213dd6f69ebce5705de5798b67c78c Mon Sep 17 00:00:00 2001
From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com>
Date: Wed, 26 Oct 2022 17:53:16 +0100
Subject: [PATCH 05/32] Add tests for cupy_ravel

---
 tests/internal/scripts/ccuda/cupy_ravel.py | 25 ++++++++++++++++++++++
 1 file changed, 25 insertions(+)
 create mode 100644 tests/internal/scripts/ccuda/cupy_ravel.py

diff --git a/tests/internal/scripts/ccuda/cupy_ravel.py b/tests/internal/scripts/ccuda/cupy_ravel.py
new file mode 100644
index 0000000000..7553f985a0
--- /dev/null
+++ b/tests/internal/scripts/ccuda/cupy_ravel.py
@@ -0,0 +1,25 @@
+from pyccel.decorators import kernel, types
+from pyccel import cuda
+import cupy as cp
+
+@kernel
+@types('int[:]')
+def func(a):
+    i = cuda.threadIdx(0) + cuda.blockIdx(0) * cuda.blockDim(0)
+    print("Hello World! ", a[i])
+
+if __name__ == '__main__':
+    threads_per_block = 5
+    n_blocks = 1
+    cp_arr = cp.array([[0, 1], [2, 3], [4, 2]])
+    flat_cp_arr = cp.array([0, 1, 2, 3, 4])
+    
+    arr1 = cp.ravel(cp_arr)
+    arr2 = cp.ravel(flat_cp_arr)
+    
+    arr2 = cp.ravel([0, 1, 2, 3, 4])
+    arr3 = cp.ravel([[0, 1], [2, 3], [4, 2]])
+    
+    cuda.deviceSynchronize()
+    func[n_blocks, threads_per_block](arr)
+    cuda.deviceSynchronize()
\ No newline at end of file

From 319d71d984329f51f614bd946ef2663913d29a3a Mon Sep 17 00:00:00 2001
From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com>
Date: Thu, 27 Oct 2022 03:55:56 +0100
Subject: [PATCH 06/32] Remove unused imports

---
 pyccel/codegen/printing/ccudacode.py | 3 +--
 1 file changed, 1 insertion(+), 2 deletions(-)

diff --git a/pyccel/codegen/printing/ccudacode.py b/pyccel/codegen/printing/ccudacode.py
index 0c087fc690..5b16bc878f 100644
--- a/pyccel/codegen/printing/ccudacode.py
+++ b/pyccel/codegen/printing/ccudacode.py
@@ -40,7 +40,7 @@
 from pyccel.ast.numpyext import NumpyFull, NumpyArray, NumpyArange
 from pyccel.ast.numpyext import NumpyReal, NumpyImag, NumpyFloat
 
-from pyccel.ast.cupyext import CupyFull, CupyArray, CupyArange, CupyRavel
+from pyccel.ast.cupyext import CupyFull, CupyArray, CupyArange
 
 from pyccel.ast.cudaext import cuda_Internal_Var, CudaArray
 
@@ -522,7 +522,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
 
-        self.add_import(c_imports['string'])
         if isinstance(arg, Variable):
             arg = self._print(arg)
             cpy_data = "cudaMemcpy({0}.raw_data, {1}.raw_data, {0}.buffer_size, cudaMemcpyHostToDevice);".format(lhs, arg)

From a1365167c04ec2c133f26d08e4f2c25a5f89c548 Mon Sep 17 00:00:00 2001
From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com>
Date: Thu, 27 Oct 2022 03:59:19 +0100
Subject: [PATCH 07/32] revert back to using more specific type for data
 pointer

---
 pyccel/codegen/printing/ccudacode.py | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/pyccel/codegen/printing/ccudacode.py b/pyccel/codegen/printing/ccudacode.py
index 5b16bc878f..ec265c5937 100644
--- a/pyccel/codegen/printing/ccudacode.py
+++ b/pyccel/codegen/printing/ccudacode.py
@@ -524,7 +524,7 @@ def copy_CudaArray_Data(self, expr):
 
         if isinstance(arg, Variable):
             arg = self._print(arg)
-            cpy_data = "cudaMemcpy({0}.raw_data, {1}.raw_data, {0}.buffer_size, cudaMemcpyHostToDevice);".format(lhs, arg)
+            cpy_data = "cudaMemcpy({0}.{2}, {1}.{2}, {0}.buffer_size, cudaMemcpyHostToDevice);".format(lhs, arg, dtype)
             return '%s\n' % (cpy_data)
         else :
             if arg.rank > 1:
@@ -532,7 +532,7 @@ def copy_CudaArray_Data(self, expr):
                 arg = self._flatten_list(arg)
             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}.{2}, {1}, {0}.buffer_size, cudaMemcpyHostToDevice);".format(self._print(lhs), dummy_array_name, dtype)
             return  '%s%s\n' % (dummy_array, cpy_data)
 
     def _print_CudaDeviceSynchronize(self, expr):

From 918a63d0c421c7f49c0c5d37f74f071d8d71ae35 Mon Sep 17 00:00:00 2001
From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com>
Date: Thu, 27 Oct 2022 14:40:20 +0100
Subject: [PATCH 08/32] Add memory location to cupy ravel

---
 pyccel/ast/cupyext.py | 15 +++++++++++----
 1 file changed, 11 insertions(+), 4 deletions(-)

diff --git a/pyccel/ast/cupyext.py b/pyccel/ast/cupyext.py
index 2c1c3f348c..5bd538f1ce 100644
--- a/pyccel/ast/cupyext.py
+++ b/pyccel/ast/cupyext.py
@@ -66,11 +66,11 @@ class CupyArray(CudaNewArray):
     arg : list, tuple, PythonList
 
     """
-    __slots__ = ('_arg','_dtype','_precision','_shape','_rank','_order')
+    __slots__ = ('_arg','_dtype','_precision','_shape','_rank','_order', '_memory_location')
     _attribute_nodes = ('_arg',)
     name = 'array'
 
-    def __init__(self, arg, dtype=None, order='C'):
+    def __init__(self, arg, dtype=None, order='C', memory_location='managed'):
 
         if not isinstance(arg, (PythonTuple, PythonList, Variable)):
             raise TypeError('Unknown type of  %s.' % type(arg))
@@ -82,6 +82,8 @@ def __init__(self, arg, dtype=None, order='C'):
         if not (is_homogeneous_tuple or is_array):
             raise TypeError('we only accept homogeneous arguments')
 
+        if memory_location not in ('host', 'device', 'managed'):
+            raise ValueError("memory_location must be 'host', 'device' or 'managed'")
         # Verify dtype and get precision
         if dtype is None:
             dtype = arg.dtype
@@ -114,6 +116,7 @@ def __init__(self, arg, dtype=None, order='C'):
         self._dtype = dtype
         self._order = order
         self._precision = prec
+        self._memory_location = memory_location
         super().__init__()
 
     def __str__(self):
@@ -123,6 +126,10 @@ def __str__(self):
     def arg(self):
         return self._arg
 
+    @property
+    def memory_location(self):
+        return self._memory_location
+
 #==============================================================================
 class CupyArange(CudaNewArray):
     """
@@ -452,13 +459,13 @@ class CupyRavel(CupyArray):
     name   = 'ravel'
     __slots__ = ('_arg','_dtype','_precision','_shape','_rank','_order', '_memory_location')
     _attribute_nodes = ('_arg',)
-    def __new__(cls, arg):
+    def __new__(cls, arg, memory_location='managed'):
         if not isinstance(arg, (list, tuple, PyccelAstNode)):
             raise TypeError('Unknown type of  %s.' % type(arg))
         return super().__new__(cls)
 
     def __init__(self, arg, memory_location='managed'):
-        super().__init__(arg)
+        super().__init__(arg = arg, memory_location = memory_location)
         shape = reduce((lambda x, y: x.python_value * y.python_value), self.shape)
         self._shape = [shape if isinstance(shape, LiteralInteger) else LiteralInteger(shape)]
         self._rank = len(self._shape)

From 1679229616016ac26711e85fa8601ac452468f1f Mon Sep 17 00:00:00 2001
From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com>
Date: Thu, 27 Oct 2022 16:30:01 +0100
Subject: [PATCH 09/32] Add cupy printer

---
 pyccel/ast/cupyext.py                |  4 ++-
 pyccel/codegen/printing/ccudacode.py | 39 ++++++++++++++++++++++++++--
 pyccel/parser/semantic.py            | 12 +++++++++
 3 files changed, 52 insertions(+), 3 deletions(-)

diff --git a/pyccel/ast/cupyext.py b/pyccel/ast/cupyext.py
index 5bd538f1ce..f2f15f122f 100644
--- a/pyccel/ast/cupyext.py
+++ b/pyccel/ast/cupyext.py
@@ -457,11 +457,13 @@ class CupyRavel(CupyArray):
             The location where the new array memory should be allocated
     """
     name   = 'ravel'
-    __slots__ = ('_arg','_dtype','_precision','_shape','_rank','_order', '_memory_location')
+    __slots__ = ('_arg','_dtype','_precision','_shape','_rank','_order', '_is_view' '_memory_location')
     _attribute_nodes = ('_arg',)
     def __new__(cls, arg, memory_location='managed'):
         if not isinstance(arg, (list, tuple, PyccelAstNode)):
             raise TypeError('Unknown type of  %s.' % type(arg))
+        if arg.rank == 0:
+            raise TypeError('Unknown type of  %s.' % type(arg))
         return super().__new__(cls)
 
     def __init__(self, arg, memory_location='managed'):
diff --git a/pyccel/codegen/printing/ccudacode.py b/pyccel/codegen/printing/ccudacode.py
index 02b0cbc069..1462846870 100644
--- a/pyccel/codegen/printing/ccudacode.py
+++ b/pyccel/codegen/printing/ccudacode.py
@@ -40,7 +40,7 @@
 from pyccel.ast.numpyext import NumpyFull, NumpyArray, NumpyArange
 from pyccel.ast.numpyext import NumpyReal, NumpyImag, NumpyFloat
 
-from pyccel.ast.cupyext import CupyFull, CupyArray, CupyArange
+from pyccel.ast.cupyext import CupyFull, CupyArray, CupyArange, CupyRavel
 
 from pyccel.ast.cudaext import cuda_Internal_Var, CudaArray
 
@@ -408,6 +408,8 @@ def _print_Assign(self, expr):
 
         if isinstance(rhs, (CupyFull)):
             return prefix_code+self.cuda_arrayFill(expr)
+        if isinstance(rhs, CupyRavel):
+            return prefix_code+self.cupy_ravel(expr)
         if isinstance(rhs, CupyArange):
             return prefix_code+self.cuda_Arange(expr)
         if isinstance(rhs, (CudaArray, CupyArray)):
@@ -497,6 +499,40 @@ def cuda_arrayFill(self, expr):
                 code_init += 'cuda_array_fill_{0}<<<1,1>>>({1}, {2});\n'.format(dtype, self._print(rhs.fill_value), self._print(lhs))
         return code_init
 
+    def cupy_ravel(self, expr):
+        """ print the assignment of a Cuda NdArray
+
+        parameters
+        ----------
+            expr : PyccelAstNode
+                The Assign Node used to get the lhs and rhs
+        Return
+        ------
+            String
+                Return a str that contains the declaration of a dummy data_buffer
+                       and a call to an operator which copies it to a Cuda NdArray struct
+                if the ndarray is a stack_array the str will contain the initialization
+        """
+        rhs = expr.rhs
+        lhs = expr.lhs
+
+        dummy_array_name = self.scope.get_new_name('cuda_array_dummy')
+        declare_dtype = self.find_in_dtype_registry(self._print(rhs.dtype), rhs.precision)
+        dtype = self.find_in_ndarray_type_registry(self._print(rhs.dtype), rhs.precision)
+        arg = rhs.arg if isinstance(rhs, (CudaArray, CupyArray)) else rhs
+
+        if isinstance(arg, Variable):
+            arg = self._print(arg)
+            cpy_data = "{0}.{2} = {1}.{2}".format(lhs, arg, dtype)
+            return '%s\n' % (cpy_data)
+        else :
+            if arg.rank > 1:
+                arg = self._flatten_list(arg)
+            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}.{2}, {1}, {0}.buffer_size, cudaMemcpyHostToDevice);".format(self._print(lhs), dummy_array_name, dtype)
+            return  '%s%s\n' % (dummy_array, cpy_data)
+
     def copy_CudaArray_Data(self, expr):
         """ print the assignment of a Cuda NdArray
 
@@ -526,7 +562,6 @@ def copy_CudaArray_Data(self, expr):
             return '%s\n' % (cpy_data)
         else :
             if arg.rank > 1:
-            # flattening the args to use them in C initialization.
                 arg = self._flatten_list(arg)
             arg = ', '.join(self._print(i) for i in arg)
             dummy_array = "%s %s[] = {%s};\n" % (declare_dtype, dummy_array_name, arg)
diff --git a/pyccel/parser/semantic.py b/pyccel/parser/semantic.py
index f2ea58c029..7009626a12 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 CupyRavel
 
 from pyccel.ast.datatypes import NativeRange, str_dtype
 from pyccel.ast.datatypes import NativeSymbol
@@ -490,6 +491,17 @@ def _infere_type(self, expr, **settings):
             d_var['cls_base'   ] = NumpyArrayClass
             return d_var
 
+        elif isinstance(expr, CupyRavel):
+            d_var['datatype'   ] = expr.dtype
+            d_var['memory_handling'] = 'heap' if not isinstance(expr.arg, Variable) 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, CudaNewArray):
             d_var['datatype'   ] = expr.dtype
             d_var['memory_handling'] = 'heap' if expr.rank > 0 else 'stack'

From dbb62a95ad5f777974f8fd2b1eb18b8b4ae5a098 Mon Sep 17 00:00:00 2001
From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com>
Date: Fri, 28 Oct 2022 12:46:36 +0100
Subject: [PATCH 10/32] Avoid calculating rank in ravel

---
 pyccel/ast/cupyext.py | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/pyccel/ast/cupyext.py b/pyccel/ast/cupyext.py
index f2f15f122f..28469cb0c8 100644
--- a/pyccel/ast/cupyext.py
+++ b/pyccel/ast/cupyext.py
@@ -470,7 +470,7 @@ def __init__(self, arg, memory_location='managed'):
         super().__init__(arg = arg, memory_location = memory_location)
         shape = reduce((lambda x, y: x.python_value * y.python_value), self.shape)
         self._shape = [shape if isinstance(shape, LiteralInteger) else LiteralInteger(shape)]
-        self._rank = len(self._shape)
+        self._rank = 1
         self._order = None
 
     @property

From dcdc81fb9230c8977b7563a658c1bedc2d1037b1 Mon Sep 17 00:00:00 2001
From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com>
Date: Fri, 4 Nov 2022 01:19:30 +0100
Subject: [PATCH 11/32] Separate tests and remove Ravel infere type

---
 pyccel/parser/semantic.py                  | 12 -----------
 tests/internal/scripts/ccuda/cupy_ravel.py | 25 ----------------------
 2 files changed, 37 deletions(-)
 delete mode 100644 tests/internal/scripts/ccuda/cupy_ravel.py

diff --git a/pyccel/parser/semantic.py b/pyccel/parser/semantic.py
index 7009626a12..f2ea58c029 100644
--- a/pyccel/parser/semantic.py
+++ b/pyccel/parser/semantic.py
@@ -61,7 +61,6 @@
 from pyccel.ast.core import Assert
 
 from pyccel.ast.class_defs import NumpyArrayClass, TupleClass, get_cls_base, CudaArrayClass
-from pyccel.ast.cupyext import CupyRavel
 
 from pyccel.ast.datatypes import NativeRange, str_dtype
 from pyccel.ast.datatypes import NativeSymbol
@@ -491,17 +490,6 @@ def _infere_type(self, expr, **settings):
             d_var['cls_base'   ] = NumpyArrayClass
             return d_var
 
-        elif isinstance(expr, CupyRavel):
-            d_var['datatype'   ] = expr.dtype
-            d_var['memory_handling'] = 'heap' if not isinstance(expr.arg, Variable) 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, CudaNewArray):
             d_var['datatype'   ] = expr.dtype
             d_var['memory_handling'] = 'heap' if expr.rank > 0 else 'stack'
diff --git a/tests/internal/scripts/ccuda/cupy_ravel.py b/tests/internal/scripts/ccuda/cupy_ravel.py
deleted file mode 100644
index 7553f985a0..0000000000
--- a/tests/internal/scripts/ccuda/cupy_ravel.py
+++ /dev/null
@@ -1,25 +0,0 @@
-from pyccel.decorators import kernel, types
-from pyccel import cuda
-import cupy as cp
-
-@kernel
-@types('int[:]')
-def func(a):
-    i = cuda.threadIdx(0) + cuda.blockIdx(0) * cuda.blockDim(0)
-    print("Hello World! ", a[i])
-
-if __name__ == '__main__':
-    threads_per_block = 5
-    n_blocks = 1
-    cp_arr = cp.array([[0, 1], [2, 3], [4, 2]])
-    flat_cp_arr = cp.array([0, 1, 2, 3, 4])
-    
-    arr1 = cp.ravel(cp_arr)
-    arr2 = cp.ravel(flat_cp_arr)
-    
-    arr2 = cp.ravel([0, 1, 2, 3, 4])
-    arr3 = cp.ravel([[0, 1], [2, 3], [4, 2]])
-    
-    cuda.deviceSynchronize()
-    func[n_blocks, threads_per_block](arr)
-    cuda.deviceSynchronize()
\ No newline at end of file

From 7c1ce50b5cf7537a90fff0bb289733f441f275f9 Mon Sep 17 00:00:00 2001
From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com>
Date: Sat, 5 Nov 2022 18:22:18 +0100
Subject: [PATCH 12/32] Implement cupy ravel with AliasAssign

---
 pyccel/codegen/printing/ccudacode.py         | 56 +++++++++++++++-----
 pyccel/parser/semantic.py                    | 15 ++++++
 pyccel/stdlib/cuda_ndarrays/cuda_ndarrays.cu | 12 +++++
 3 files changed, 70 insertions(+), 13 deletions(-)

diff --git a/pyccel/codegen/printing/ccudacode.py b/pyccel/codegen/printing/ccudacode.py
index 1462846870..c9e910f362 100644
--- a/pyccel/codegen/printing/ccudacode.py
+++ b/pyccel/codegen/printing/ccudacode.py
@@ -408,10 +408,10 @@ def _print_Assign(self, expr):
 
         if isinstance(rhs, (CupyFull)):
             return prefix_code+self.cuda_arrayFill(expr)
-        if isinstance(rhs, CupyRavel):
-            return prefix_code+self.cupy_ravel(expr)
         if isinstance(rhs, CupyArange):
             return prefix_code+self.cuda_Arange(expr)
+        if isinstance(rhs, CupyRavel):
+            return prefix_code+self.cupy_ravel(expr)
         if isinstance(rhs, (CudaArray, CupyArray)):
             return prefix_code+self.copy_CudaArray_Data(expr)
         if isinstance(rhs, (NumpyArray, PythonTuple)):
@@ -424,6 +424,41 @@ def _print_Assign(self, expr):
         rhs = self._print(expr.rhs)
         return prefix_code+'{} = {};\n'.format(lhs, rhs)
 
+    def _print_AliasAssign(self, expr):
+        lhs_var = expr.lhs
+        rhs_var = expr.rhs
+
+        lhs_address = ObjectAddress(lhs_var)
+        rhs_address = ObjectAddress(rhs_var)
+
+        # the below condition handles the case of reassinging a pointer to an array view.
+        # setting the pointer's is_view attribute to false so it can be ignored by the free_pointer function.
+
+        if not self.stored_in_c_pointer(lhs_var) and \
+                isinstance(lhs_var, Variable) and lhs_var.is_ndarray:
+            # rhs = self._print(rhs_var)
+            if isinstance(rhs_var, CupyRavel):
+                memory_location = rhs_var.memory_location
+                if memory_location in ('device', 'host'):
+                    memory_location = 'allocateMemoryOn' + str(memory_location).capitalize()
+                else:
+                    memory_location = 'managedMemory'
+                return 'cupy_ravel({}, {}, {});\n'.format(lhs_var, rhs_var, memory_location)
+            elif isinstance(rhs_var, Variable) and rhs_var.is_ndarray:
+                lhs = self._print(lhs_address)
+                if lhs_var.order == rhs_var.order:
+                    return 'alias_assign({}, {});\n'.format(lhs, rhs)
+                else:
+                    return 'transpose_alias_assign({}, {});\n'.format(lhs, rhs)
+            else:
+                lhs = self._print(lhs_var)
+                return '{} = {};\n'.format(lhs, rhs)
+        else:
+            lhs = self._print(lhs_address)
+            rhs = self._print(rhs_address)
+
+            return '{} = {};\n'.format(lhs, rhs)
+
     def arrayFill(self, expr):
         """ print the assignment of a NdArray
 
@@ -521,17 +556,12 @@ def cupy_ravel(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 isinstance(arg, Variable):
-            arg = self._print(arg)
-            cpy_data = "{0}.{2} = {1}.{2}".format(lhs, arg, dtype)
-            return '%s\n' % (cpy_data)
-        else :
-            if arg.rank > 1:
-                arg = self._flatten_list(arg)
-            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}.{2}, {1}, {0}.buffer_size, cudaMemcpyHostToDevice);".format(self._print(lhs), dummy_array_name, dtype)
-            return  '%s%s\n' % (dummy_array, cpy_data)
+        if arg.rank > 1:
+            arg = self._flatten_list(arg)
+        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}.{2}, {1}, {0}.buffer_size, cudaMemcpyHostToDevice);".format(self._print(lhs), dummy_array_name, dtype)
+        return  '%s%s\n' % (dummy_array, cpy_data)
 
     def copy_CudaArray_Data(self, expr):
         """ print the assignment of a Cuda NdArray
diff --git a/pyccel/parser/semantic.py b/pyccel/parser/semantic.py
index f2ea58c029..1d569c85e2 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 CupyRavel
 
 from pyccel.ast.datatypes import NativeRange, str_dtype
 from pyccel.ast.datatypes import NativeSymbol
@@ -490,6 +491,20 @@ def _infere_type(self, expr, **settings):
             d_var['cls_base'   ] = NumpyArrayClass
             return d_var
 
+        elif isinstance(expr, CupyRavel):
+
+            var = expr.arg
+
+            d_var['memory_handling'] = 'alias' if isinstance(var, Variable) else 'heap'
+            d_var['memory_location'] = expr.memory_location
+            d_var['datatype'      ] = var.dtype
+            d_var['shape'         ] = tuple(reversed(var.shape))
+            d_var['rank'          ] = var.rank
+            d_var['cls_base'      ] = CudaNewArray
+            d_var['order'         ] = var.order
+            d_var['precision'     ] = var.precision
+            return d_var
+
         elif isinstance(expr, CudaNewArray):
             d_var['datatype'   ] = expr.dtype
             d_var['memory_handling'] = 'heap' if expr.rank > 0 else 'stack'
diff --git a/pyccel/stdlib/cuda_ndarrays/cuda_ndarrays.cu b/pyccel/stdlib/cuda_ndarrays/cuda_ndarrays.cu
index 2c0d517e19..10674ffbb8 100644
--- a/pyccel/stdlib/cuda_ndarrays/cuda_ndarrays.cu
+++ b/pyccel/stdlib/cuda_ndarrays/cuda_ndarrays.cu
@@ -132,6 +132,18 @@ t_ndarray   cuda_array_create(int32_t nd, int64_t *shape,
     return (arr);
 }
 
+void cupy_ravel(t_ndarray *dest, t_ndarray2 src, enum e_memory_locations location)
+{
+    *dest = src;
+    void (*fun_ptr_arr[])(void*, size_t) = {managed_memory, host_memory, device_memory};
+
+    (*fun_ptr_arr[location])(&(dest.shape), sizeof(int64_t));
+    (*fun_ptr_arr[location])(&(dest.strides), sizeof(int64_t));
+    *(dest->shape) = src.lenght;
+    *(dest->strides) = 1;
+    dest->is_view = true;
+}
+
 int32_t cuda_free_array(t_ndarray arr)
 {
     if (arr.shape == NULL)

From adcf8649471ffe79da698d77d96cb8557f35cfa8 Mon Sep 17 00:00:00 2001
From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com>
Date: Sat, 5 Nov 2022 19:18:00 +0100
Subject: [PATCH 13/32] Fix address printing in cupy_ravel

---
 pyccel/codegen/printing/ccudacode.py | 8 +++++---
 1 file changed, 5 insertions(+), 3 deletions(-)

diff --git a/pyccel/codegen/printing/ccudacode.py b/pyccel/codegen/printing/ccudacode.py
index c9e910f362..900ac66c55 100644
--- a/pyccel/codegen/printing/ccudacode.py
+++ b/pyccel/codegen/printing/ccudacode.py
@@ -436,15 +436,17 @@ def _print_AliasAssign(self, expr):
 
         if not self.stored_in_c_pointer(lhs_var) and \
                 isinstance(lhs_var, Variable) and lhs_var.is_ndarray:
-            # rhs = self._print(rhs_var)
             if isinstance(rhs_var, CupyRavel):
+                rhs = self._print(rhs_var.arg)
                 memory_location = rhs_var.memory_location
                 if memory_location in ('device', 'host'):
                     memory_location = 'allocateMemoryOn' + str(memory_location).capitalize()
                 else:
                     memory_location = 'managedMemory'
-                return 'cupy_ravel({}, {}, {});\n'.format(lhs_var, rhs_var, memory_location)
-            elif isinstance(rhs_var, Variable) and rhs_var.is_ndarray:
+                lhs = self._print(lhs_address)
+                return 'cupy_ravel({}, {}, {});\n'.format(lhs, rhs, memory_location)
+            rhs = self._print(rhs_var)
+            if isinstance(rhs_var, Variable) and rhs_var.is_ndarray:
                 lhs = self._print(lhs_address)
                 if lhs_var.order == rhs_var.order:
                     return 'alias_assign({}, {});\n'.format(lhs, rhs)

From d93856122e445e129c3a2717c31f99f24b5a90c9 Mon Sep 17 00:00:00 2001
From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com>
Date: Sat, 5 Nov 2022 23:24:46 +0100
Subject: [PATCH 14/32] Fix typo in cuda_ravel implementation

---
 pyccel/stdlib/cuda_ndarrays/cuda_ndarrays.cu | 5 +++--
 1 file changed, 3 insertions(+), 2 deletions(-)

diff --git a/pyccel/stdlib/cuda_ndarrays/cuda_ndarrays.cu b/pyccel/stdlib/cuda_ndarrays/cuda_ndarrays.cu
index 10674ffbb8..143b9ad163 100644
--- a/pyccel/stdlib/cuda_ndarrays/cuda_ndarrays.cu
+++ b/pyccel/stdlib/cuda_ndarrays/cuda_ndarrays.cu
@@ -137,8 +137,9 @@ void cupy_ravel(t_ndarray *dest, t_ndarray2 src, enum e_memory_locations locatio
     *dest = src;
     void (*fun_ptr_arr[])(void*, size_t) = {managed_memory, host_memory, device_memory};
 
-    (*fun_ptr_arr[location])(&(dest.shape), sizeof(int64_t));
-    (*fun_ptr_arr[location])(&(dest.strides), sizeof(int64_t));
+    dest->nd = nd;
+    (*fun_ptr_arr[location])(&(dest->shape), sizeof(int64_t));
+    (*fun_ptr_arr[location])(&(dest->strides), sizeof(int64_t));
     *(dest->shape) = src.lenght;
     *(dest->strides) = 1;
     dest->is_view = true;

From a16bbeefb4e6e7e9b038b868cd8f4df00e5ff594 Mon Sep 17 00:00:00 2001
From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com>
Date: Sat, 5 Nov 2022 23:39:24 +0100
Subject: [PATCH 15/32] Fix case of passed variable

---
 pyccel/codegen/printing/ccudacode.py | 31 ----------------------------
 pyccel/parser/semantic.py            | 15 ++++++--------
 2 files changed, 6 insertions(+), 40 deletions(-)

diff --git a/pyccel/codegen/printing/ccudacode.py b/pyccel/codegen/printing/ccudacode.py
index 900ac66c55..8de59d27f5 100644
--- a/pyccel/codegen/printing/ccudacode.py
+++ b/pyccel/codegen/printing/ccudacode.py
@@ -410,8 +410,6 @@ def _print_Assign(self, expr):
             return prefix_code+self.cuda_arrayFill(expr)
         if isinstance(rhs, CupyArange):
             return prefix_code+self.cuda_Arange(expr)
-        if isinstance(rhs, CupyRavel):
-            return prefix_code+self.cupy_ravel(expr)
         if isinstance(rhs, (CudaArray, CupyArray)):
             return prefix_code+self.copy_CudaArray_Data(expr)
         if isinstance(rhs, (NumpyArray, PythonTuple)):
@@ -536,35 +534,6 @@ def cuda_arrayFill(self, expr):
                 code_init += 'cuda_array_fill_{0}<<<1,1>>>({1}, {2});\n'.format(dtype, self._print(rhs.fill_value), self._print(lhs))
         return code_init
 
-    def cupy_ravel(self, expr):
-        """ print the assignment of a Cuda NdArray
-
-        parameters
-        ----------
-            expr : PyccelAstNode
-                The Assign Node used to get the lhs and rhs
-        Return
-        ------
-            String
-                Return a str that contains the declaration of a dummy data_buffer
-                       and a call to an operator which copies it to a Cuda NdArray struct
-                if the ndarray is a stack_array the str will contain the initialization
-        """
-        rhs = expr.rhs
-        lhs = expr.lhs
-
-        dummy_array_name = self.scope.get_new_name('cuda_array_dummy')
-        declare_dtype = self.find_in_dtype_registry(self._print(rhs.dtype), rhs.precision)
-        dtype = self.find_in_ndarray_type_registry(self._print(rhs.dtype), rhs.precision)
-        arg = rhs.arg if isinstance(rhs, (CudaArray, CupyArray)) else rhs
-
-        if arg.rank > 1:
-            arg = self._flatten_list(arg)
-        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}.{2}, {1}, {0}.buffer_size, cudaMemcpyHostToDevice);".format(self._print(lhs), dummy_array_name, dtype)
-        return  '%s%s\n' % (dummy_array, cpy_data)
-
     def copy_CudaArray_Data(self, expr):
         """ print the assignment of a Cuda NdArray
 
diff --git a/pyccel/parser/semantic.py b/pyccel/parser/semantic.py
index 1d569c85e2..57aa87cf62 100644
--- a/pyccel/parser/semantic.py
+++ b/pyccel/parser/semantic.py
@@ -492,17 +492,14 @@ def _infere_type(self, expr, **settings):
             return d_var
 
         elif isinstance(expr, CupyRavel):
-
-            var = expr.arg
-
-            d_var['memory_handling'] = 'alias' if isinstance(var, Variable) else 'heap'
+            d_var['memory_handling'] = 'alias' if isinstance(expr.arg, Variable) else 'heap'
             d_var['memory_location'] = expr.memory_location
-            d_var['datatype'      ] = var.dtype
-            d_var['shape'         ] = tuple(reversed(var.shape))
-            d_var['rank'          ] = var.rank
+            d_var['datatype'      ] = expr.dtype
+            d_var['shape'         ] = expr.shape
+            d_var['rank'          ] = expr.rank
+            d_var['precision'     ] = expr.precision
+            d_var['order'         ] = None
             d_var['cls_base'      ] = CudaNewArray
-            d_var['order'         ] = var.order
-            d_var['precision'     ] = var.precision
             return d_var
 
         elif isinstance(expr, CudaNewArray):

From 587844666993b3f1bb5c6719745dcd1f4dc148e5 Mon Sep 17 00:00:00 2001
From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com>
Date: Sun, 6 Nov 2022 00:29:26 +0100
Subject: [PATCH 16/32] avoid call to _print

---
 pyccel/codegen/printing/ccudacode.py |  3 +--
 pyccel/parser/semantic.py            | 12 ++++++------
 2 files changed, 7 insertions(+), 8 deletions(-)

diff --git a/pyccel/codegen/printing/ccudacode.py b/pyccel/codegen/printing/ccudacode.py
index 8de59d27f5..09757c8554 100644
--- a/pyccel/codegen/printing/ccudacode.py
+++ b/pyccel/codegen/printing/ccudacode.py
@@ -435,14 +435,13 @@ def _print_AliasAssign(self, expr):
         if not self.stored_in_c_pointer(lhs_var) and \
                 isinstance(lhs_var, Variable) and lhs_var.is_ndarray:
             if isinstance(rhs_var, CupyRavel):
-                rhs = self._print(rhs_var.arg)
                 memory_location = rhs_var.memory_location
                 if memory_location in ('device', 'host'):
                     memory_location = 'allocateMemoryOn' + str(memory_location).capitalize()
                 else:
                     memory_location = 'managedMemory'
                 lhs = self._print(lhs_address)
-                return 'cupy_ravel({}, {}, {});\n'.format(lhs, rhs, memory_location)
+                return 'cupy_ravel({}, {}, {});\n'.format(lhs, rhs_var, memory_location)
             rhs = self._print(rhs_var)
             if isinstance(rhs_var, Variable) and rhs_var.is_ndarray:
                 lhs = self._print(lhs_address)
diff --git a/pyccel/parser/semantic.py b/pyccel/parser/semantic.py
index 57aa87cf62..089b36022b 100644
--- a/pyccel/parser/semantic.py
+++ b/pyccel/parser/semantic.py
@@ -494,12 +494,12 @@ def _infere_type(self, expr, **settings):
         elif isinstance(expr, CupyRavel):
             d_var['memory_handling'] = 'alias' if isinstance(expr.arg, Variable) else 'heap'
             d_var['memory_location'] = expr.memory_location
-            d_var['datatype'      ] = expr.dtype
-            d_var['shape'         ] = expr.shape
-            d_var['rank'          ] = expr.rank
-            d_var['precision'     ] = expr.precision
-            d_var['order'         ] = None
-            d_var['cls_base'      ] = CudaNewArray
+            d_var['datatype'       ] = expr.dtype
+            d_var['shape'          ] = expr.shape
+            d_var['rank'           ] = expr.rank
+            d_var['precision'      ] = expr.precision
+            d_var['order'          ] = None
+            d_var['cls_base'       ] = CudaNewArray
             return d_var
 
         elif isinstance(expr, CudaNewArray):

From a895e7141640cf4a9ec2552be0a2e77710139772 Mon Sep 17 00:00:00 2001
From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com>
Date: Sun, 6 Nov 2022 15:16:38 +0100
Subject: [PATCH 17/32] Fix typo in dest nd

---
 pyccel/stdlib/cuda_ndarrays/cuda_ndarrays.cu | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/pyccel/stdlib/cuda_ndarrays/cuda_ndarrays.cu b/pyccel/stdlib/cuda_ndarrays/cuda_ndarrays.cu
index 143b9ad163..9d929d9ee9 100644
--- a/pyccel/stdlib/cuda_ndarrays/cuda_ndarrays.cu
+++ b/pyccel/stdlib/cuda_ndarrays/cuda_ndarrays.cu
@@ -137,7 +137,7 @@ void cupy_ravel(t_ndarray *dest, t_ndarray2 src, enum e_memory_locations locatio
     *dest = src;
     void (*fun_ptr_arr[])(void*, size_t) = {managed_memory, host_memory, device_memory};
 
-    dest->nd = nd;
+    dest->nd = 1;
     (*fun_ptr_arr[location])(&(dest->shape), sizeof(int64_t));
     (*fun_ptr_arr[location])(&(dest->strides), sizeof(int64_t));
     *(dest->shape) = src.lenght;

From 0c9c07082ad181282efd071f75b49ac1fada564c Mon Sep 17 00:00:00 2001
From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com>
Date: Fri, 11 Nov 2022 16:21:25 +0100
Subject: [PATCH 18/32] Remove memory location option for cupy functions

---
 pyccel/ast/cupyext.py | 15 ++++++---------
 1 file changed, 6 insertions(+), 9 deletions(-)

diff --git a/pyccel/ast/cupyext.py b/pyccel/ast/cupyext.py
index 28469cb0c8..eaefdec731 100644
--- a/pyccel/ast/cupyext.py
+++ b/pyccel/ast/cupyext.py
@@ -70,7 +70,7 @@ class CupyArray(CudaNewArray):
     _attribute_nodes = ('_arg',)
     name = 'array'
 
-    def __init__(self, arg, dtype=None, order='C', memory_location='managed'):
+    def __init__(self, arg, dtype=None, order='C'):
 
         if not isinstance(arg, (PythonTuple, PythonList, Variable)):
             raise TypeError('Unknown type of  %s.' % type(arg))
@@ -82,8 +82,6 @@ def __init__(self, arg, dtype=None, order='C', memory_location='managed'):
         if not (is_homogeneous_tuple or is_array):
             raise TypeError('we only accept homogeneous arguments')
 
-        if memory_location not in ('host', 'device', 'managed'):
-            raise ValueError("memory_location must be 'host', 'device' or 'managed'")
         # Verify dtype and get precision
         if dtype is None:
             dtype = arg.dtype
@@ -116,7 +114,7 @@ def __init__(self, arg, dtype=None, order='C', memory_location='managed'):
         self._dtype = dtype
         self._order = order
         self._precision = prec
-        self._memory_location = memory_location
+        self._memory_location = 'device'
         super().__init__()
 
     def __str__(self):
@@ -457,21 +455,20 @@ class CupyRavel(CupyArray):
             The location where the new array memory should be allocated
     """
     name   = 'ravel'
-    __slots__ = ('_arg','_dtype','_precision','_shape','_rank','_order', '_is_view' '_memory_location')
+    __slots__ = ('_arg','_dtype','_precision','_shape','_rank','_order', '_memory_location')
     _attribute_nodes = ('_arg',)
-    def __new__(cls, arg, memory_location='managed'):
+    def __new__(cls, arg):
         if not isinstance(arg, (list, tuple, PyccelAstNode)):
             raise TypeError('Unknown type of  %s.' % type(arg))
         if arg.rank == 0:
             raise TypeError('Unknown type of  %s.' % type(arg))
         return super().__new__(cls)
 
-    def __init__(self, arg, memory_location='managed'):
-        super().__init__(arg = arg, memory_location = memory_location)
+    def __init__(self, arg):
+        super().__init__(arg = arg)
         shape = reduce((lambda x, y: x.python_value * y.python_value), self.shape)
         self._shape = [shape if isinstance(shape, LiteralInteger) else LiteralInteger(shape)]
         self._rank = 1
-        self._order = None
 
     @property
     def arg(self):

From 573e03bc7ea641667e20624062f660abfadb935f Mon Sep 17 00:00:00 2001
From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com>
Date: Fri, 11 Nov 2022 16:22:22 +0100
Subject: [PATCH 19/32] Avoid if statement for memory location

---
 pyccel/codegen/printing/ccudacode.py | 6 +-----
 1 file changed, 1 insertion(+), 5 deletions(-)

diff --git a/pyccel/codegen/printing/ccudacode.py b/pyccel/codegen/printing/ccudacode.py
index 09757c8554..41f734ffd6 100644
--- a/pyccel/codegen/printing/ccudacode.py
+++ b/pyccel/codegen/printing/ccudacode.py
@@ -435,11 +435,7 @@ def _print_AliasAssign(self, expr):
         if not self.stored_in_c_pointer(lhs_var) and \
                 isinstance(lhs_var, Variable) and lhs_var.is_ndarray:
             if isinstance(rhs_var, CupyRavel):
-                memory_location = rhs_var.memory_location
-                if memory_location in ('device', 'host'):
-                    memory_location = 'allocateMemoryOn' + str(memory_location).capitalize()
-                else:
-                    memory_location = 'managedMemory'
+                memory_location = 'allocateMemoryOnDevice'
                 lhs = self._print(lhs_address)
                 return 'cupy_ravel({}, {}, {});\n'.format(lhs, rhs_var, memory_location)
             rhs = self._print(rhs_var)

From 3d97f1725cab365d8c2d1e8cd54c98b7e471beef Mon Sep 17 00:00:00 2001
From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com>
Date: Fri, 11 Nov 2022 16:23:11 +0100
Subject: [PATCH 20/32] Add tests for cupy ravel

---
 tests/internal/scripts/ccuda/cupy_ravel.py    | 18 +++++++++++++++++
 .../scripts/ccuda/cupy_ravel_variable.py      | 20 +++++++++++++++++++
 2 files changed, 38 insertions(+)
 create mode 100644 tests/internal/scripts/ccuda/cupy_ravel.py
 create mode 100644 tests/internal/scripts/ccuda/cupy_ravel_variable.py

diff --git a/tests/internal/scripts/ccuda/cupy_ravel.py b/tests/internal/scripts/ccuda/cupy_ravel.py
new file mode 100644
index 0000000000..7afb31f9e3
--- /dev/null
+++ b/tests/internal/scripts/ccuda/cupy_ravel.py
@@ -0,0 +1,18 @@
+from pyccel.decorators import kernel, types
+from pyccel import cuda
+import cupy as cp
+
+@kernel
+@types('int[:]', 'int[:]')
+def func(a, b):
+    i = cuda.threadIdx(0) + cuda.blockIdx(0) * cuda.blockDim(0)
+    print("Hello World! ", a[i], b[i])
+
+if __name__ == '__main__':
+    threads_per_block = 32
+    n_blocks = 1
+    arr1 = cp.ravel([[1,2],[1,3]])
+    arr2 = cp.ravel([1,2,3,4])
+    cuda.deviceSynchronize()
+    func[n_blocks, threads_per_block](arr1, arr2)
+    cuda.deviceSynchronize()
\ No newline at end of file
diff --git a/tests/internal/scripts/ccuda/cupy_ravel_variable.py b/tests/internal/scripts/ccuda/cupy_ravel_variable.py
new file mode 100644
index 0000000000..2165ff6e62
--- /dev/null
+++ b/tests/internal/scripts/ccuda/cupy_ravel_variable.py
@@ -0,0 +1,20 @@
+from pyccel.decorators import kernel, types
+from pyccel import cuda
+import cupy as cp
+
+@kernel
+@types('int[:]', 'int[:]')
+def func(a, b):
+    i = cuda.threadIdx(0) + cuda.blockIdx(0) * cuda.blockDim(0)
+    print("Hello World! ", a[i], b[i])
+
+if __name__ == '__main__':
+    threads_per_block = 32
+    n_blocks = 1
+    c =[[1,2],[1,3]]
+    arr = cuda.array(c, dtype=int)
+    arr1 = cp.ravel(arr)
+    arr2 = cp.ravel(c)
+    cuda.deviceSynchronize()
+    func[n_blocks, threads_per_block](arr1, arr2)
+    cuda.deviceSynchronize()
\ No newline at end of file

From 6127038cc896d988b4609ca51a46bcf2d5c17dd3 Mon Sep 17 00:00:00 2001
From: EmilyBoune <louise.bourne@gmail.com>
Date: Fri, 30 Dec 2022 12:07:24 +0100
Subject: [PATCH 21/32] Line endings

---
 pyccel/ast/cudaext.py                               | 2 +-
 tests/internal/scripts/ccuda/cupy_ravel.py          | 2 +-
 tests/internal/scripts/ccuda/cupy_ravel_variable.py | 2 +-
 3 files changed, 3 insertions(+), 3 deletions(-)

diff --git a/pyccel/ast/cudaext.py b/pyccel/ast/cudaext.py
index aed2e206ed..e34d325b73 100644
--- a/pyccel/ast/cudaext.py
+++ b/pyccel/ast/cudaext.py
@@ -198,4 +198,4 @@ class CudaGridDim(CudaInternalVar)   : pass
 }
 cuda_mod = Module('cuda',
     variables = cuda_constants.values(),
-    funcs = cuda_funcs.values())
\ No newline at end of file
+    funcs = cuda_funcs.values())
diff --git a/tests/internal/scripts/ccuda/cupy_ravel.py b/tests/internal/scripts/ccuda/cupy_ravel.py
index 7afb31f9e3..c9b0525839 100644
--- a/tests/internal/scripts/ccuda/cupy_ravel.py
+++ b/tests/internal/scripts/ccuda/cupy_ravel.py
@@ -15,4 +15,4 @@ def func(a, b):
     arr2 = cp.ravel([1,2,3,4])
     cuda.deviceSynchronize()
     func[n_blocks, threads_per_block](arr1, arr2)
-    cuda.deviceSynchronize()
\ No newline at end of file
+    cuda.deviceSynchronize()
diff --git a/tests/internal/scripts/ccuda/cupy_ravel_variable.py b/tests/internal/scripts/ccuda/cupy_ravel_variable.py
index 2165ff6e62..b144834a00 100644
--- a/tests/internal/scripts/ccuda/cupy_ravel_variable.py
+++ b/tests/internal/scripts/ccuda/cupy_ravel_variable.py
@@ -17,4 +17,4 @@ def func(a, b):
     arr2 = cp.ravel(c)
     cuda.deviceSynchronize()
     func[n_blocks, threads_per_block](arr1, arr2)
-    cuda.deviceSynchronize()
\ No newline at end of file
+    cuda.deviceSynchronize()

From 2a94f14930c44510f04fecc5a98c2b3f627642b4 Mon Sep 17 00:00:00 2001
From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com>
Date: Fri, 30 Dec 2022 20:58:36 +0100
Subject: [PATCH 22/32] Remove slot from CupyRavel

Co-authored-by: EmilyBourne <louise.bourne@gmail.com>
---
 pyccel/ast/cupyext.py | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/pyccel/ast/cupyext.py b/pyccel/ast/cupyext.py
index eaefdec731..69bbe61121 100644
--- a/pyccel/ast/cupyext.py
+++ b/pyccel/ast/cupyext.py
@@ -455,7 +455,7 @@ class CupyRavel(CupyArray):
             The location where the new array memory should be allocated
     """
     name   = 'ravel'
-    __slots__ = ('_arg','_dtype','_precision','_shape','_rank','_order', '_memory_location')
+    __slots__ = ()
     _attribute_nodes = ('_arg',)
     def __new__(cls, arg):
         if not isinstance(arg, (list, tuple, PyccelAstNode)):

From 27c42509cd49ed0d43dc62a1788e1bec1b4d5c12 Mon Sep 17 00:00:00 2001
From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com>
Date: Sat, 31 Dec 2022 19:22:54 +0100
Subject: [PATCH 23/32] Fix typos in cupy_ravel

---
 pyccel/stdlib/cuda_ndarrays/cuda_ndarrays.cu | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/pyccel/stdlib/cuda_ndarrays/cuda_ndarrays.cu b/pyccel/stdlib/cuda_ndarrays/cuda_ndarrays.cu
index 9d929d9ee9..d708d746d7 100644
--- a/pyccel/stdlib/cuda_ndarrays/cuda_ndarrays.cu
+++ b/pyccel/stdlib/cuda_ndarrays/cuda_ndarrays.cu
@@ -132,7 +132,7 @@ t_ndarray   cuda_array_create(int32_t nd, int64_t *shape,
     return (arr);
 }
 
-void cupy_ravel(t_ndarray *dest, t_ndarray2 src, enum e_memory_locations location)
+void cupy_ravel(t_ndarray *dest, t_ndarray src, enum e_memory_locations location)
 {
     *dest = src;
     void (*fun_ptr_arr[])(void*, size_t) = {managed_memory, host_memory, device_memory};
@@ -140,7 +140,7 @@ void cupy_ravel(t_ndarray *dest, t_ndarray2 src, enum e_memory_locations locatio
     dest->nd = 1;
     (*fun_ptr_arr[location])(&(dest->shape), sizeof(int64_t));
     (*fun_ptr_arr[location])(&(dest->strides), sizeof(int64_t));
-    *(dest->shape) = src.lenght;
+    *(dest->shape) = src.length;
     *(dest->strides) = 1;
     dest->is_view = true;
 }

From d83f52634a287a4fd29620abdfece94ea0a29720 Mon Sep 17 00:00:00 2001
From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com>
Date: Sun, 1 Jan 2023 15:59:41 +0100
Subject: [PATCH 24/32] Fix codacy problems

---
 pyccel/ast/cupyext.py                             |  2 ++
 pyccel/codegen/printing/ccudacode.py              | 15 ++++++++-------
 tests/internal/scripts/ccuda/cupy_ravel.py        |  8 +++++---
 .../internal/scripts/ccuda/cupy_ravel_variable.py |  8 +++++---
 4 files changed, 20 insertions(+), 13 deletions(-)

diff --git a/pyccel/ast/cupyext.py b/pyccel/ast/cupyext.py
index 69bbe61121..1c56d7da52 100644
--- a/pyccel/ast/cupyext.py
+++ b/pyccel/ast/cupyext.py
@@ -126,6 +126,8 @@ def arg(self):
 
     @property
     def memory_location(self):
+        """ Indicate if the array is in the host or device memory
+        """
         return self._memory_location
 
 #==============================================================================
diff --git a/pyccel/codegen/printing/ccudacode.py b/pyccel/codegen/printing/ccudacode.py
index aae499f96c..b1edbbfa45 100644
--- a/pyccel/codegen/printing/ccudacode.py
+++ b/pyccel/codegen/printing/ccudacode.py
@@ -439,22 +439,22 @@ def _print_AliasAssign(self, expr):
             if isinstance(rhs_var, CupyRavel):
                 memory_location = 'allocateMemoryOnDevice'
                 lhs = self._print(lhs_address)
-                return 'cupy_ravel({}, {}, {});\n'.format(lhs, rhs_var, memory_location)
+                return f'cupy_ravel({lhs}, {rhs_var}, {memory_location});\n'
             rhs = self._print(rhs_var)
             if isinstance(rhs_var, Variable) and rhs_var.is_ndarray:
                 lhs = self._print(lhs_address)
                 if lhs_var.order == rhs_var.order:
-                    return 'alias_assign({}, {});\n'.format(lhs, rhs)
+                    return f'alias_assign({lhs}, {rhs});\n'
                 else:
-                    return 'transpose_alias_assign({}, {});\n'.format(lhs, rhs)
+                    return f'transpose_alias_assign({lhs}, {rhs});\n'
             else:
                 lhs = self._print(lhs_var)
-                return '{} = {};\n'.format(lhs, rhs)
+                return f'{lhs} = {rhs};\n'
         else:
             lhs = self._print(lhs_address)
             rhs = self._print(rhs_address)
 
-            return '{} = {};\n'.format(lhs, rhs)
+            return f'{lhs} = {rhs};\n'
 
     def arrayFill(self, expr):
         """ print the assignment of a NdArray
@@ -556,14 +556,15 @@ def copy_CudaArray_Data(self, expr):
 
         if isinstance(arg, Variable):
             arg = self._print(arg)
-            cpy_data = "cudaMemcpy({0}.{2}, {1}.{2}, {0}.buffer_size, cudaMemcpyHostToDevice);".format(lhs, arg, dtype)
+            cpy_data = f"cudaMemcpy({lhs}.{dtype}, {arg}.{dtype}, {lhs}.buffer_size, cudaMemcpyHostToDevice);"
             return '%s\n' % (cpy_data)
         else :
             if arg.rank > 1:
                 arg = self._flatten_list(arg)
             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}.{2}, {1}, {0}.buffer_size, cudaMemcpyHostToDevice);".format(self._print(lhs), dummy_array_name, dtype)
+            target_array_name = self._print(lhs)
+            cpy_data = f"cudaMemcpy({target_array_name}.{dtype}, {dummy_array_name}, {target_array_name}.buffer_size, cudaMemcpyHostToDevice);"
             return  '%s%s\n' % (dummy_array, cpy_data)
 
     def _print_CudaSynchronize(self, expr):
diff --git a/tests/internal/scripts/ccuda/cupy_ravel.py b/tests/internal/scripts/ccuda/cupy_ravel.py
index c9b0525839..0d5ad20786 100644
--- a/tests/internal/scripts/ccuda/cupy_ravel.py
+++ b/tests/internal/scripts/ccuda/cupy_ravel.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[:]', 'int[:]')
@@ -13,6 +15,6 @@ def func(a, b):
     n_blocks = 1
     arr1 = cp.ravel([[1,2],[1,3]])
     arr2 = cp.ravel([1,2,3,4])
-    cuda.deviceSynchronize()
+    cuda.synchronize()
     func[n_blocks, threads_per_block](arr1, arr2)
-    cuda.deviceSynchronize()
+    cuda.synchronize()
diff --git a/tests/internal/scripts/ccuda/cupy_ravel_variable.py b/tests/internal/scripts/ccuda/cupy_ravel_variable.py
index b144834a00..832625260f 100644
--- a/tests/internal/scripts/ccuda/cupy_ravel_variable.py
+++ b/tests/internal/scripts/ccuda/cupy_ravel_variable.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[:]', 'int[:]')
@@ -15,6 +17,6 @@ def func(a, b):
     arr = cuda.array(c, dtype=int)
     arr1 = cp.ravel(arr)
     arr2 = cp.ravel(c)
-    cuda.deviceSynchronize()
+    cuda.synchronize()
     func[n_blocks, threads_per_block](arr1, arr2)
-    cuda.deviceSynchronize()
+    cuda.synchronize()

From 5ff5c22f840cb2033f9cb71db26a402eb4d1e9f9 Mon Sep 17 00:00:00 2001
From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com>
Date: Mon, 2 Jan 2023 20:06:24 +0100
Subject: [PATCH 25/32] Use managed_memory to allocate shape and strides

---
 pyccel/codegen/printing/ccudacode.py         | 3 +--
 pyccel/stdlib/cuda_ndarrays/cuda_ndarrays.cu | 8 +++-----
 2 files changed, 4 insertions(+), 7 deletions(-)

diff --git a/pyccel/codegen/printing/ccudacode.py b/pyccel/codegen/printing/ccudacode.py
index b1edbbfa45..128db6e81e 100644
--- a/pyccel/codegen/printing/ccudacode.py
+++ b/pyccel/codegen/printing/ccudacode.py
@@ -437,9 +437,8 @@ def _print_AliasAssign(self, expr):
         if not self.stored_in_c_pointer(lhs_var) and \
                 isinstance(lhs_var, Variable) and lhs_var.is_ndarray:
             if isinstance(rhs_var, CupyRavel):
-                memory_location = 'allocateMemoryOnDevice'
                 lhs = self._print(lhs_address)
-                return f'cupy_ravel({lhs}, {rhs_var}, {memory_location});\n'
+                return f'cupy_ravel({lhs}, {rhs_var});\n'
             rhs = self._print(rhs_var)
             if isinstance(rhs_var, Variable) and rhs_var.is_ndarray:
                 lhs = self._print(lhs_address)
diff --git a/pyccel/stdlib/cuda_ndarrays/cuda_ndarrays.cu b/pyccel/stdlib/cuda_ndarrays/cuda_ndarrays.cu
index d708d746d7..ab600f2c25 100644
--- a/pyccel/stdlib/cuda_ndarrays/cuda_ndarrays.cu
+++ b/pyccel/stdlib/cuda_ndarrays/cuda_ndarrays.cu
@@ -132,14 +132,12 @@ t_ndarray   cuda_array_create(int32_t nd, int64_t *shape,
     return (arr);
 }
 
-void cupy_ravel(t_ndarray *dest, t_ndarray src, enum e_memory_locations location)
+void cupy_ravel(t_ndarray *dest, t_ndarray src)
 {
     *dest = src;
-    void (*fun_ptr_arr[])(void*, size_t) = {managed_memory, host_memory, device_memory};
-
     dest->nd = 1;
-    (*fun_ptr_arr[location])(&(dest->shape), sizeof(int64_t));
-    (*fun_ptr_arr[location])(&(dest->strides), sizeof(int64_t));
+    managed_memory(&(dest->shape), sizeof(int64_t));
+    managed_memory(&(dest->strides), sizeof(int64_t));
     *(dest->shape) = src.length;
     *(dest->strides) = 1;
     dest->is_view = true;

From d472a7cf0454f499371d53ddb50dee180ca2d313 Mon Sep 17 00:00:00 2001
From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com>
Date: Fri, 6 Jan 2023 11:50:23 +0100
Subject: [PATCH 26/32] Do a copy for host arrays

---
 pyccel/parser/semantic.py                        |  5 ++++-
 pyccel/stdlib/cuda_ndarrays/cuda_ndarrays.h      |  4 +++-
 .../scripts/ccuda/cupy_ravel_variable.py         | 16 +++++++++-------
 3 files changed, 16 insertions(+), 9 deletions(-)

diff --git a/pyccel/parser/semantic.py b/pyccel/parser/semantic.py
index 089b36022b..cf43d36d13 100644
--- a/pyccel/parser/semantic.py
+++ b/pyccel/parser/semantic.py
@@ -492,7 +492,10 @@ def _infere_type(self, expr, **settings):
             return d_var
 
         elif isinstance(expr, CupyRavel):
-            d_var['memory_handling'] = 'alias' if isinstance(expr.arg, Variable) else 'heap'
+            if isinstance(expr.arg, Variable) and expr.arg.memory_location == "device":
+                d_var['memory_handling'] = 'alias'
+            else:
+                d_var['memory_handling'] = 'heap'
             d_var['memory_location'] = expr.memory_location
             d_var['datatype'       ] = expr.dtype
             d_var['shape'          ] = expr.shape
diff --git a/pyccel/stdlib/cuda_ndarrays/cuda_ndarrays.h b/pyccel/stdlib/cuda_ndarrays/cuda_ndarrays.h
index d0cce3b5f0..c928d859e7 100644
--- a/pyccel/stdlib/cuda_ndarrays/cuda_ndarrays.h
+++ b/pyccel/stdlib/cuda_ndarrays/cuda_ndarrays.h
@@ -21,7 +21,9 @@ void _cuda_array_fill_int64(int64_t c, t_ndarray arr);
 __global__
 void _cuda_array_fill_double(double c, t_ndarray arr);
 
-t_ndarray   cuda_array_create(int32_t nd, int64_t *shape, enum e_types type, bool is_view, enum e_memory_locations location);
+void            cupy_ravel(t_ndarray *dest, t_ndarray src);
+
+t_ndarray       cuda_array_create(int32_t nd, int64_t *shape, enum e_types type, bool is_view, enum e_memory_locations location);
 int32_t         cuda_free_array(t_ndarray dump);
 int32_t         cuda_free_pointer(t_ndarray dump);
 #endif
diff --git a/tests/internal/scripts/ccuda/cupy_ravel_variable.py b/tests/internal/scripts/ccuda/cupy_ravel_variable.py
index 832625260f..65550ae6d3 100644
--- a/tests/internal/scripts/ccuda/cupy_ravel_variable.py
+++ b/tests/internal/scripts/ccuda/cupy_ravel_variable.py
@@ -5,18 +5,20 @@
 from pyccel import cuda
 
 @kernel
-@types('int[:]', 'int[:]')
-def func(a, b):
+@types('int[:]', 'int[:]', 'int[:]')
+def func(a, b, c):
     i = cuda.threadIdx(0) + cuda.blockIdx(0) * cuda.blockDim(0)
-    print("Hello World! ", a[i], b[i])
+    print("Hello World! ", a[i], b[i], c[i])
 
 if __name__ == '__main__':
     threads_per_block = 32
     n_blocks = 1
     c =[[1,2],[1,3]]
-    arr = cuda.array(c, dtype=int)
-    arr1 = cp.ravel(arr)
-    arr2 = cp.ravel(c)
+    host_arr = cuda.array(c, dtype=int)
+    device_arr = cuda.array(c, dtype=int, memory_location='device')
+    arr1 = cp.ravel(host_arr)
+    arr2 = cp.ravel(device_arr)
+    arr3 = cp.ravel(c)
     cuda.synchronize()
-    func[n_blocks, threads_per_block](arr1, arr2)
+    func[n_blocks, threads_per_block](arr1, arr2, arr3)
     cuda.synchronize()

From be2fd091c234f25b4ceaddf4fb8ac361e6c16966 Mon Sep 17 00:00:00 2001
From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com>
Date: Sat, 7 Jan 2023 04:58:26 +0100
Subject: [PATCH 27/32] Remove _attribute_nodes from CupyRavel

---
 pyccel/ast/cupyext.py | 1 -
 1 file changed, 1 deletion(-)

diff --git a/pyccel/ast/cupyext.py b/pyccel/ast/cupyext.py
index 1c56d7da52..f9684e3b58 100644
--- a/pyccel/ast/cupyext.py
+++ b/pyccel/ast/cupyext.py
@@ -458,7 +458,6 @@ class CupyRavel(CupyArray):
     """
     name   = 'ravel'
     __slots__ = ()
-    _attribute_nodes = ('_arg',)
     def __new__(cls, arg):
         if not isinstance(arg, (list, tuple, PyccelAstNode)):
             raise TypeError('Unknown type of  %s.' % type(arg))

From 3fb754ffb4b4eb90b92fcb423f01a41c6843fa50 Mon Sep 17 00:00:00 2001
From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com>
Date: Sat, 7 Jan 2023 05:05:46 +0100
Subject: [PATCH 28/32] Use tuple instead of list for variable in test

---
 tests/internal/scripts/ccuda/cupy_ravel_variable.py | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/tests/internal/scripts/ccuda/cupy_ravel_variable.py b/tests/internal/scripts/ccuda/cupy_ravel_variable.py
index 65550ae6d3..5dcc5e42d5 100644
--- a/tests/internal/scripts/ccuda/cupy_ravel_variable.py
+++ b/tests/internal/scripts/ccuda/cupy_ravel_variable.py
@@ -13,7 +13,7 @@ def func(a, b, c):
 if __name__ == '__main__':
     threads_per_block = 32
     n_blocks = 1
-    c =[[1,2],[1,3]]
+    c = ((1, 2), (1, 3))
     host_arr = cuda.array(c, dtype=int)
     device_arr = cuda.array(c, dtype=int, memory_location='device')
     arr1 = cp.ravel(host_arr)

From def5e4eb1e30a14de57c60d1033c70f4d0561c20 Mon Sep 17 00:00:00 2001
From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com>
Date: Sat, 7 Jan 2023 14:59:37 +0100
Subject: [PATCH 29/32] Remove unnecessary arg property getter CupyRavel

---
 pyccel/ast/cupyext.py | 3 ---
 1 file changed, 3 deletions(-)

diff --git a/pyccel/ast/cupyext.py b/pyccel/ast/cupyext.py
index f9684e3b58..04e46b8be9 100644
--- a/pyccel/ast/cupyext.py
+++ b/pyccel/ast/cupyext.py
@@ -471,9 +471,6 @@ def __init__(self, arg):
         self._shape = [shape if isinstance(shape, LiteralInteger) else LiteralInteger(shape)]
         self._rank = 1
 
-    @property
-    def arg(self):
-        return self._arg
 
 #==============================================================================
 cupy_funcs = {

From 6d8489d76c704f09456a8e9664996c16d00309c9 Mon Sep 17 00:00:00 2001
From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com>
Date: Tue, 10 Jan 2023 02:15:36 +0100
Subject: [PATCH 30/32] Add case of PyccelArraySize for cupy ravel

---
 pyccel/ast/cupyext.py | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/pyccel/ast/cupyext.py b/pyccel/ast/cupyext.py
index 04e46b8be9..9c413d5bd5 100644
--- a/pyccel/ast/cupyext.py
+++ b/pyccel/ast/cupyext.py
@@ -468,7 +468,7 @@ def __new__(cls, arg):
     def __init__(self, arg):
         super().__init__(arg = arg)
         shape = reduce((lambda x, y: x.python_value * y.python_value), self.shape)
-        self._shape = [shape if isinstance(shape, LiteralInteger) else LiteralInteger(shape)]
+        self._shape = [shape if isinstance(shape, (LiteralInteger, PyccelArraySize)) else LiteralInteger(shape)]
         self._rank = 1
 
 

From c970bf7c2a765b2460d73a623965a1bd1f052511 Mon Sep 17 00:00:00 2001
From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com>
Date: Wed, 11 Jan 2023 01:25:35 +0100
Subject: [PATCH 31/32] Use cudaMallocManaged in cupy_ravel

---
 pyccel/stdlib/cuda_ndarrays/cuda_ndarrays.cu | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/pyccel/stdlib/cuda_ndarrays/cuda_ndarrays.cu b/pyccel/stdlib/cuda_ndarrays/cuda_ndarrays.cu
index ab600f2c25..e284971c35 100644
--- a/pyccel/stdlib/cuda_ndarrays/cuda_ndarrays.cu
+++ b/pyccel/stdlib/cuda_ndarrays/cuda_ndarrays.cu
@@ -136,8 +136,8 @@ void cupy_ravel(t_ndarray *dest, t_ndarray src)
 {
     *dest = src;
     dest->nd = 1;
-    managed_memory(&(dest->shape), sizeof(int64_t));
-    managed_memory(&(dest->strides), sizeof(int64_t));
+    cudaMallocManaged(&(dest->shape), sizeof(int64_t));
+    cudaMallocManaged(&(dest->strides), sizeof(int64_t));
     *(dest->shape) = src.length;
     *(dest->strides) = 1;
     dest->is_view = true;

From a8c5db4779b6f2030eb1da75ce1191a1ee3df333 Mon Sep 17 00:00:00 2001
From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com>
Date: Wed, 11 Jan 2023 01:26:06 +0100
Subject: [PATCH 32/32] Clean up cupy ravel tests

---
 tests/internal/scripts/ccuda/cupy_ravel.py          |  9 ---------
 tests/internal/scripts/ccuda/cupy_ravel_variable.py | 12 ------------
 2 files changed, 21 deletions(-)

diff --git a/tests/internal/scripts/ccuda/cupy_ravel.py b/tests/internal/scripts/ccuda/cupy_ravel.py
index 0d5ad20786..97e3342df1 100644
--- a/tests/internal/scripts/ccuda/cupy_ravel.py
+++ b/tests/internal/scripts/ccuda/cupy_ravel.py
@@ -4,17 +4,8 @@
 from pyccel.decorators import kernel, types
 from pyccel import cuda
 
-@kernel
-@types('int[:]', 'int[:]')
-def func(a, b):
-    i = cuda.threadIdx(0) + cuda.blockIdx(0) * cuda.blockDim(0)
-    print("Hello World! ", a[i], b[i])
-
 if __name__ == '__main__':
     threads_per_block = 32
     n_blocks = 1
     arr1 = cp.ravel([[1,2],[1,3]])
     arr2 = cp.ravel([1,2,3,4])
-    cuda.synchronize()
-    func[n_blocks, threads_per_block](arr1, arr2)
-    cuda.synchronize()
diff --git a/tests/internal/scripts/ccuda/cupy_ravel_variable.py b/tests/internal/scripts/ccuda/cupy_ravel_variable.py
index 5dcc5e42d5..1062c262ba 100644
--- a/tests/internal/scripts/ccuda/cupy_ravel_variable.py
+++ b/tests/internal/scripts/ccuda/cupy_ravel_variable.py
@@ -1,24 +1,12 @@
 # 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
 
-@kernel
-@types('int[:]', 'int[:]', 'int[:]')
-def func(a, b, c):
-    i = cuda.threadIdx(0) + cuda.blockIdx(0) * cuda.blockDim(0)
-    print("Hello World! ", a[i], b[i], c[i])
-
 if __name__ == '__main__':
-    threads_per_block = 32
-    n_blocks = 1
     c = ((1, 2), (1, 3))
     host_arr = cuda.array(c, dtype=int)
     device_arr = cuda.array(c, dtype=int, memory_location='device')
     arr1 = cp.ravel(host_arr)
     arr2 = cp.ravel(device_arr)
     arr3 = cp.ravel(c)
-    cuda.synchronize()
-    func[n_blocks, threads_per_block](arr1, arr2, arr3)
-    cuda.synchronize()