cuda.host_empty function #67

smazouz42 · 2024-07-22T15:31:11Z

This pull request addresses issue #56 by adding a new feature to 'cuda' host_empty that allows you to allocate memory on the CPU

This PR aims to make the C code compilable using nvcc. The cuda language was added as well as a CudaCodePrinter. Changes to stdlib: Wrapped expressions using complex types in an `ifndef __NVCC__` to avoid processing them with the nvcc compiler --------- Co-authored-by: Mouad Elalj, EmilyBourne

This pull request fixes #48, by implementing a tiny wrapper for CUDA and a wrapper for non-CUDA functionalities only with external 'C'. **Commit Summary** - Implemented new header printer for CUDA. - Added CUDA wrapper assignment - Instead of wrapping all local headers, wrap only C functions with extern 'C' --------- Co-authored-by: EmilyBourne <[email protected]> Co-authored-by: bauom <[email protected]>

This pull request addresses issue #28 by implementing a new feature in Pyccel that allows users to define custom GPU kernels. The syntax for creating these kernels is inspired by Numba. and I also need to fix issue #45 for testing purposes **Commit Summary** - Introduced KernelCall class - Added cuda printer methods _print_KernelCall and _print_FunctionDef to generate the corresponding CUDA representation for both kernel calls and definitions - Added IndexedFunctionCall represents an indexed function call - Added CUDA module and cuda.synchronize() - Fixing a bug that I found in the header: it does not import the necessary header for the used function --------- Co-authored-by: EmilyBourne <[email protected]> Co-authored-by: bauom <[email protected]> Co-authored-by: Emily Bourne <[email protected]>

…nctions, and refining CUDA type handling

This PR aims to make the C code compilable using nvcc. The cuda language was added as well as a CudaCodePrinter. Changes to stdlib: Wrapped expressions using complex types in an `ifndef __NVCC__` to avoid processing them with the nvcc compiler --------- Co-authored-by: Mouad Elalj, EmilyBourne

This pull request fixes #48, by implementing a tiny wrapper for CUDA and a wrapper for non-CUDA functionalities only with external 'C'. **Commit Summary** - Implemented new header printer for CUDA. - Added CUDA wrapper assignment - Instead of wrapping all local headers, wrap only C functions with extern 'C' --------- Co-authored-by: EmilyBourne <[email protected]> Co-authored-by: bauom <[email protected]>

This pull request addresses issue #28 by implementing a new feature in Pyccel that allows users to define custom GPU kernels. The syntax for creating these kernels is inspired by Numba. and I also need to fix issue #45 for testing purposes **Commit Summary** - Introduced KernelCall class - Added cuda printer methods _print_KernelCall and _print_FunctionDef to generate the corresponding CUDA representation for both kernel calls and definitions - Added IndexedFunctionCall represents an indexed function call - Added CUDA module and cuda.synchronize() - Fixing a bug that I found in the header: it does not import the necessary header for the used function --------- Co-authored-by: EmilyBourne <[email protected]> Co-authored-by: bauom <[email protected]> Co-authored-by: Emily Bourne <[email protected]>

This pull request addresses issue #59 by adding more CUDA-specific keywords to enhance the checking of variable/function names and prevent name clashes --------- Co-authored-by: EmilyBourne <[email protected]> Co-authored-by: bauom <[email protected]>

This pull request addresses issue #41 by implementing a new feature in Pyccel that allows users to define a custom device **Commit Summary** - Adding handler for custom device and its code generation. - Adding test --------- Co-authored-by: EmilyBourne <[email protected]>

pyccel-bot · 2024-07-24T21:15:01Z

Unfortunately your PR is not passing the tests so it is not quite ready for review yet. Let me know when it is fixed with /bot mark as ready.

pyccel-bot

There seems to be lines in this PR which aren't tested. Please take a look at my comments and add tests which cover the new code.

If this is modified code which cannot be easily tested in this PR please open an issue to request that this code be either removed or tested. Once you have done that please leave a message on the relevant conversation beginning with the line /bot accept and referencing the issue.

Similarly if the new code cannot be tested for some reason, please leave a comment beginning with the line /bot accept on the relevant conversation explaining why the code can't be tested.

pyccel-bot · 2024-07-24T21:32:24Z

pyccel/ast/cudatypes.py

+            return NotImplemented
+


This code isn't tested. Please can you take a look

/bot accept
The fallback return NotImplemented does not need testing

smazouz42 · 2024-07-25T09:56:36Z

/bot run docs

pyccel-bot

Good job ! Your PR is using all the code it added/changed.

pyccel-bot · 2024-07-25T10:20:56Z

@jalalium, @smazouz42 has been working hard and thinks that they have now replied to or fixed all your comments. Could you take another look at the PR and see if you can approve now?

This PR aims to make the C code compilable using nvcc. The cuda language was added as well as a CudaCodePrinter. Changes to stdlib: Wrapped expressions using complex types in an `ifndef __NVCC__` to avoid processing them with the nvcc compiler --------- Co-authored-by: Mouad Elalj, EmilyBourne

This pull request fixes #48, by implementing a tiny wrapper for CUDA and a wrapper for non-CUDA functionalities only with external 'C'. **Commit Summary** - Implemented new header printer for CUDA. - Added CUDA wrapper assignment - Instead of wrapping all local headers, wrap only C functions with extern 'C' --------- Co-authored-by: EmilyBourne <[email protected]> Co-authored-by: bauom <[email protected]>

This pull request addresses issue #28 by implementing a new feature in Pyccel that allows users to define custom GPU kernels. The syntax for creating these kernels is inspired by Numba. and I also need to fix issue #45 for testing purposes **Commit Summary** - Introduced KernelCall class - Added cuda printer methods _print_KernelCall and _print_FunctionDef to generate the corresponding CUDA representation for both kernel calls and definitions - Added IndexedFunctionCall represents an indexed function call - Added CUDA module and cuda.synchronize() - Fixing a bug that I found in the header: it does not import the necessary header for the used function --------- Co-authored-by: EmilyBourne <[email protected]> Co-authored-by: bauom <[email protected]> Co-authored-by: Emily Bourne <[email protected]>

This pull request addresses issue #59 by adding more CUDA-specific keywords to enhance the checking of variable/function names and prevent name clashes --------- Co-authored-by: EmilyBourne <[email protected]> Co-authored-by: bauom <[email protected]>

This pull request addresses issue #41 by implementing a new feature in Pyccel that allows users to define a custom device **Commit Summary** - Adding handler for custom device and its code generation. - Adding test --------- Co-authored-by: EmilyBourne <[email protected]>

…sue_56

jalalium

Very good job!

pyccel-bot · 2024-07-26T15:51:08Z

Hey @yguclu, @EmilyBourne, this PR is looking pretty good. @smazouz42 and @jalalium think it is ready to merge. Could you add your expertise to confirm that this follows all the coding conventions and fits in Pyccel's future plans? Thanks 😄

EmilyBourne · 2024-07-29T16:07:42Z

docs/cuda.md

+
+### cuda+host_empty
+
+The cuda+host_empty function allocates an empty array on the host.


Is cuda+host_empty cuda.host_empty?

EmilyBourne · 2024-07-29T16:08:03Z

pyccel/ast/class_defs.py

+    'IntegerClass',
+    'FloatClass',


These classes now appear twice in this list

EmilyBourne · 2024-07-29T16:08:15Z

pyccel/ast/class_defs.py

    elif isinstance(class_type, (NumpyNumericType, NumpyNDArrayType)):
        return NumpyArrayClass
+    # elif isinstance(class_type, StackArrayType):


Suggested change

# elif isinstance(class_type, StackArrayType):

EmilyBourne · 2024-07-29T16:09:42Z

pyccel/ast/cudaext.py

+    def __init__(self, *args ,class_type, init_dtype, memory_location):
+        self._class_type = class_type
+        self._init_dtype = init_dtype
+        self._memory_location = memory_location


Is the memory location not inside the class type already?

Suggested change

def __init__(self, *args ,class_type, init_dtype, memory_location):

self._class_type = class_type

self._init_dtype = init_dtype

self._memory_location = memory_location

def __init__(self, *args, class_type, init_dtype, memory_location):

self._class_type = class_type

self._init_dtype = init_dtype

self._memory_location = memory_location

EmilyBourne · 2024-07-29T16:12:48Z

pyccel/ast/cudaext.py

+    'full'              : PyccelFunctionDef('full' , CudaFull),
+    'host_empty'             : PyccelFunctionDef('host_empty' , CudaHostEmpty),


Suggested change

'full' : PyccelFunctionDef('full' , CudaFull),

'host_empty' : PyccelFunctionDef('host_empty' , CudaHostEmpty),

'full' : PyccelFunctionDef('full' , CudaFull),

'host_empty' : PyccelFunctionDef('host_empty' , CudaHostEmpty),

EmilyBourne · 2024-07-29T16:23:03Z

pyccel/codegen/printing/cucode.py

+        if isinstance(rhs.class_type, CudaArrayType):
+            if(isinstance(rhs, (CudaFull))):
+            # TODO add support for CudaFull
+                return " \n"


Is it not safer to not include this code so that the neat error is raised instead of just printing nothing?

EmilyBourne · 2024-07-29T16:23:18Z

pyccel/codegen/utilities.py

+                                             accelerators=('python',))),
+    "numpy_f90"    : ("numpy", CompileObj("numpy_f90.f90",folder="numpy")),
+    "numpy_c"      : ("numpy", CompileObj("numpy_c.c",folder="numpy")),
+    "Set_extensions" : ("STC_Extensions", CompileObj("Set_Extensions.h",


Why this change? Bad merge?

EmilyBourne · 2024-07-29T16:24:23Z

pyccel/cuda/cuda_arrays.py

+    array
+        The empty array on the host.
+    """
+    import numpy as np


Usually in Python best practice is to place all imports at the top of the file. Is there a reason you don't do that here?

EmilyBourne · 2024-07-29T16:25:08Z

pyccel/stdlib/cuda_ndarrays/cuda_ndarrays.cu

+    return (1);
+}
+
+__host__ __device__
+int32_t cuda_free(t_ndarray  arr)
+{
+    if (arr.shape == NULL)
+        return (0);
+    cudaFree(arr.raw_data);
+    arr.raw_data = NULL;
+    cudaFree(arr.shape);
+    arr.shape = NULL;
+    cudaFree(arr.strides);
+    arr.strides = NULL;
+    return (0);


What does the 1/0 returned for host/device freeing represent?

EmilyBourne · 2024-07-29T16:26:13Z

pyproject.toml

@@ -58,7 +58,8 @@ include = [
  "pyccel/stdlib/**/*.c",
  "pyccel/stdlib/**/*.f90",
  "pyccel/extensions/STC/include",
-  "pyccel/extensions/gFTL/include/v2"
+  "pyccel/extensions/gFTL/include/v2",
+  "pyccel/stdlib/cuda_ndarrays/cuda_ndarrays.cu"


Please group this with the other stdlib files. I think we can safely include all cuda files found inside stdlib.

EmilyBourne and others added 30 commits June 27, 2024 08:10

Trigger tests on push to devel or main branch

c7a6638

Add cuda workflow to test cuda developments on CI

821a1c5

Trigger tests on push to devel or main branch

092b557

Begin implementation of CUDA arrays: adding cudaempty and cudafull fu…

80f905b

…nctions, and refining CUDA type handling

work in progress

7e8cf9e

work in progress

2dbcfae

work in progress

f3911d5

work in progress

37289f9

work in progress

ba66b48

work in progress

406a88b

work in progress

3afad1b

work in progress

190c5a2

cleaning up my PR

eeeb249

cleaning up my PR

de0f5ab

cleaning up my PR

d6ba6ad

work in progress

8286a89

work in progress

96c3f29

work in progress

b414d62

Trigger tests on push to devel or main branch

7c93416

Add cuda workflow to test cuda developments on CI

f8ec722

Trigger tests on push to devel or main branch

cc3a93e

work in progress

a28c724

pyccel-bot bot suggested changes Jul 24, 2024

View reviewed changes

fix doc string of host_empty

eea028a

smazouz42 marked this pull request as ready for review July 25, 2024 09:59

pyccel-bot bot approved these changes Jul 25, 2024

View reviewed changes

pyccel-bot bot added the needs_initial_review label Jul 25, 2024

pyccel-bot bot requested a review from jalalium July 25, 2024 10:20

EmilyBourne and others added 8 commits July 26, 2024 14:08

Trigger tests on push to devel or main branch

cc5a8cf

Add cuda workflow to test cuda developments on CI

a822c41

Trigger tests on push to devel or main branch

99b1838

EmilyBourne force-pushed the devel branch from 8eef19d to 12d98b6 Compare July 26, 2024 12:09

Merge branch 'devel' of https://github.com/pyccel/pyccel-cuda into is…

09c6b74

…sue_56

jalalium approved these changes Jul 26, 2024

View reviewed changes

pyccel-bot bot added Ready_for_review Received at least one approval. Requires review from senior developer and removed needs_initial_review labels Jul 26, 2024

pyccel-bot bot requested review from EmilyBourne and yguclu July 26, 2024 15:51

EmilyBourne requested changes Jul 29, 2024

View reviewed changes

EmilyBourne force-pushed the devel branch 2 times, most recently from 81b9970 to 5f7e3e2 Compare September 3, 2024 13:43

EmilyBourne force-pushed the devel branch from 5f7e3e2 to bb18b0a Compare September 25, 2024 15:40

EmilyBourne force-pushed the devel branch from bb18b0a to de362d3 Compare November 8, 2024 13:54

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

cuda.host_empty function #67

cuda.host_empty function #67

smazouz42 commented Jul 22, 2024

pyccel-bot bot commented Jul 24, 2024

pyccel-bot bot left a comment

pyccel-bot bot Jul 24, 2024

smazouz42 Jul 25, 2024

smazouz42 commented Jul 25, 2024

pyccel-bot bot left a comment

pyccel-bot bot commented Jul 25, 2024

jalalium left a comment

pyccel-bot bot commented Jul 26, 2024

EmilyBourne Jul 29, 2024

EmilyBourne Jul 29, 2024

EmilyBourne Jul 29, 2024

EmilyBourne Jul 29, 2024

EmilyBourne Jul 29, 2024

EmilyBourne Jul 29, 2024

EmilyBourne Jul 29, 2024

EmilyBourne Jul 29, 2024

EmilyBourne Jul 29, 2024

EmilyBourne Jul 29, 2024


		### cuda+host_empty

		The cuda+host_empty function allocates an empty array on the host.

		'full' : PyccelFunctionDef('full' , CudaFull),
		'host_empty' : PyccelFunctionDef('host_empty' , CudaHostEmpty),

cuda.host_empty function #67

Are you sure you want to change the base?

cuda.host_empty function #67

Conversation

smazouz42 commented Jul 22, 2024

pyccel-bot bot commented Jul 24, 2024

pyccel-bot bot left a comment

Choose a reason for hiding this comment

pyccel-bot bot Jul 24, 2024

Choose a reason for hiding this comment

Choose a reason for hiding this comment

smazouz42 commented Jul 25, 2024

pyccel-bot bot left a comment

Choose a reason for hiding this comment

pyccel-bot bot commented Jul 25, 2024

jalalium left a comment

Choose a reason for hiding this comment

pyccel-bot bot commented Jul 26, 2024

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment