work in progress

Author: smazouz42

pyccel/ast/cudatypes.py

@@ -32,11 +32,6 @@ class CudaArrayType(HomogeneousContainerType, metaclass = ArgumentSingleton):
     """
     __slots__ = ('_element_type', '_container_rank', '_order', '_memory_location')
 
-    # def __new__(cls, dtype, rank, order, memory_location):
-    #     if rank == 0:
-    #         return dtype
-    #     else:
-    #         return super().__new__(cls, dtype, rank, order)
     def __init__(self, dtype, rank, order, memory_location):
         assert isinstance(rank, int)
         assert order in (None, 'C', 'F')

pyccel/codegen/printing/ccode.py

@@ -1315,8 +1315,7 @@ def get_declare_type(self, expr):
                 self.add_import(c_imports['ndarrays'])
                 dtype = 't_ndarray'
             elif isinstance(expr.class_type, CudaArrayType):
-                self.add_import(c_imports['ndarrays'])
-                dtype = 't_ndarray'
+                dtype = 't_cuda_ndarray'
 
             else:
                 errors.report(PYCCEL_RESTRICTION_TODO+' (rank>0)', symbol=expr, severity='fatal')

pyccel/codegen/printing/cucode.py

@@ -163,7 +163,6 @@ def _print_Allocate(self, expr):
         else:
             memory_location = 'managedMemory'
         self.add_import(c_imports['cuda_ndarrays'])
-        self.add_import(c_imports['ndarrays'])
         alloc_code = f"{self._print(expr.variable)} = cuda_array_create({variable.rank},  shape_Assign, {dtype}, {is_view},{memory_location});\n"
         return f'{shape_Assign} {alloc_code}'
 

pyccel/stdlib/cuda_ndarrays/cuda_ndarrays.cu

@@ -14,10 +14,10 @@ void    host_memory(void** devPtr, size_t size)
 {
     cudaMallocHost(devPtr, size);
 }
-t_ndarray   cuda_array_create(enum e_memory_locations location, int32_t nd, int64_t *shape,
+t_cuda_ndarray   cuda_array_create(enum e_memory_locations location, int32_t nd, int64_t *shape,
         enum e_types type, bool is_view)
 {
-    t_ndarray arr;
+    t_cuda_ndarray arr;
     void (*fun_ptr_arr[])(void**, size_t) = {managed_memory, host_memory, device_memory};
 
     arr.nd = nd;
@@ -61,7 +61,7 @@ t_ndarray   cuda_array_create(enum e_memory_locations location, int32_t nd, int6
     return (arr);
 }
 
-int32_t cuda_free_host(t_ndarray arr)
+int32_t cuda_free_host(t_cuda_ndarray arr)
 {
     if (arr.shape == NULL)
         return (0);
@@ -75,7 +75,7 @@ int32_t cuda_free_host(t_ndarray arr)
 }
 
 __host__ __device__
-int32_t cuda_free(t_ndarray arr)
+int32_t cuda_free(t_cuda_ndarray arr)
 {
     if (arr.shape == NULL)
         return (0);
@@ -87,7 +87,7 @@ int32_t cuda_free(t_ndarray arr)
 }
 
 __host__ __device__
-int32_t cuda_free_pointer(t_ndarray arr)
+int32_t cuda_free_pointer(t_cuda_ndarray arr)
 {
     if (arr.is_view == false || arr.shape == NULL)
         return (0);

pyccel/stdlib/cuda_ndarrays/cuda_ndarrays.h

@@ -3,11 +3,62 @@
 
 # include <cuda_runtime.h>
 # include <iostream>
-#include "../ndarrays/ndarrays.h"
 
-t_ndarray   cuda_array_create(int32_t nd, int64_t *shape, enum e_types type, bool is_view ,
+typedef enum e_types
+{
+        nd_bool     = 0,
+        nd_int8     = 1,
+        nd_int16    = 3,
+        nd_int32    = 5,
+        nd_int64    = 7,
+        nd_float    = 11,
+        nd_double   = 12,
+        nd_cfloat   = 14,
+        nd_cdouble  = 15
+} t_types;
+
+
+enum e_memory_locations
+{
+        managedMemory,
+        allocateMemoryOnHost,
+        allocateMemoryOnDevice
+};
+
+typedef enum e_order
+{
+    order_f,
+    order_c,
+} t_order;
+
+typedef struct  s_cuda_ndarray
+{
+    void            *raw_data;
+    /* number of dimensions */
+    int32_t                 nd;
+    /* shape 'size of each dimension' */
+    int64_t                 *shape;
+    /* strides 'number of elements to skip to get the next element' */
+    int64_t                 *strides;
+    /* type of the array elements */
+    t_types            type;
+    /* type size of the array elements */
+    int32_t                 type_size;
+    /* number of element in the array */
+    int32_t                 length;
+    /* size of the array */
+    int32_t                 buffer_size;
+    /* True if the array does not own the data */
+    bool                    is_view;
+    /* stores the order of the array: order_f or order_c */
+    t_order            order;
+}               t_cuda_ndarray;
+
+
+t_cuda_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_host(t_ndarray arr);
+int32_t cuda_free_host(t_cuda_ndarray arr);
+
 
 
 using namespace std;