From 71c37d693bbff99bfb105256e446552695da22dc Mon Sep 17 00:00:00 2001 From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com> Date: Fri, 18 Nov 2022 12:31:53 +0100 Subject: [PATCH 1/6] Make CudaDeviceSynchronize a void function --- pyccel/ast/cudaext.py | 6 +++--- pyccel/codegen/printing/ccode.py | 3 ++- pyccel/codegen/printing/ccudacode.py | 2 +- pyccel/parser/semantic.py | 6 ++++-- tests/internal/scripts/ccuda/cuda_array_managed.py | 2 +- 5 files changed, 11 insertions(+), 8 deletions(-) diff --git a/pyccel/ast/cudaext.py b/pyccel/ast/cudaext.py index 79f42bda74..3c6c6ca992 100644 --- a/pyccel/ast/cudaext.py +++ b/pyccel/ast/cudaext.py @@ -17,7 +17,7 @@ from .datatypes import (dtype_and_precision_registry as dtype_registry, default_precision, datatype, NativeInteger, - NativeFloat, NativeComplex, NativeBool, str_dtype, + NativeFloat, NativeComplex, NativeBool, NativeVoid, str_dtype, NativeNumeric) from .internals import PyccelInternalFunction, Slice, max_precision, get_final_precision @@ -145,8 +145,8 @@ def __init__(self): #... self._shape = None self._rank = 0 - self._dtype = NativeInteger() - self._precision = None + self._dtype = NativeVoid() + self._precision = 0 self._order = None super().__init__() diff --git a/pyccel/codegen/printing/ccode.py b/pyccel/codegen/printing/ccode.py index 6300ffb982..0e80fbcf35 100644 --- a/pyccel/codegen/printing/ccode.py +++ b/pyccel/codegen/printing/ccode.py @@ -209,7 +209,8 @@ ('int',8) : 'int64_t', ('int',2) : 'int16_t', ('int',1) : 'int8_t', - ('bool',4) : 'bool'} + ('bool',4) : 'bool', + ('void',0) : 'void',} ndarray_type_registry = { ('float',8) : 'nd_double', diff --git a/pyccel/codegen/printing/ccudacode.py b/pyccel/codegen/printing/ccudacode.py index 7bbf6f0d67..191d6ccdb2 100644 --- a/pyccel/codegen/printing/ccudacode.py +++ b/pyccel/codegen/printing/ccudacode.py @@ -535,7 +535,7 @@ def copy_CudaArray_Data(self, expr): return '%s%s\n' % (dummy_array, cpy_data) def _print_CudaDeviceSynchronize(self, expr): - return 'cudaDeviceSynchronize()' + return 'cudaDeviceSynchronize();\n' def _print_CudaInternalVar(self, expr): var_name = type(expr).__name__ diff --git a/pyccel/parser/semantic.py b/pyccel/parser/semantic.py index 53c022993f..eadbb427cc 100644 --- a/pyccel/parser/semantic.py +++ b/pyccel/parser/semantic.py @@ -67,7 +67,7 @@ from pyccel.ast.datatypes import default_precision from pyccel.ast.datatypes import (NativeInteger, NativeBool, NativeFloat, NativeString, - NativeGeneric, NativeComplex) + NativeGeneric, NativeComplex, NativeVoid) from pyccel.ast.functionalexpr import FunctionalSum, FunctionalMax, FunctionalMin, GeneratorComprehension, FunctionalFor @@ -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') @@ -2574,6 +2574,8 @@ def _visit_Assign(self, expr, **settings): bounding_box=(self._current_fst_node.lineno, self._current_fst_node.col_offset), severity='error') return None + if lhs.is_temp and rhs.dtype is NativeVoid(): + return rhs lhs = self._assign_lhs_variable(lhs, d_var, rhs, new_expressions, isinstance(expr, AugAssign), **settings) elif isinstance(lhs, PythonTuple): n = len(lhs) diff --git a/tests/internal/scripts/ccuda/cuda_array_managed.py b/tests/internal/scripts/ccuda/cuda_array_managed.py index 5d55151e65..4e53545929 100644 --- a/tests/internal/scripts/ccuda/cuda_array_managed.py +++ b/tests/internal/scripts/ccuda/cuda_array_managed.py @@ -12,6 +12,6 @@ def square(a): n_blocks = 1 a = cuda.array([0,1,2,3,4], memory_location = 'managed') cuda.deviceSynchronize() - square[n_blocks, threads_per_block](a) + square[5, 1](a) cuda.deviceSynchronize() print(a) \ No newline at end of file From 42c994203c1136fb2642cb9972c317d421389383 Mon Sep 17 00:00:00 2001 From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com> Date: Fri, 30 Dec 2022 21:34:48 +0100 Subject: [PATCH 2/6] Only Apply NativeVoid check on PyccelSymbol --- 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 eadbb427cc..02bf689b69 100644 --- a/pyccel/parser/semantic.py +++ b/pyccel/parser/semantic.py @@ -2574,7 +2574,7 @@ def _visit_Assign(self, expr, **settings): bounding_box=(self._current_fst_node.lineno, self._current_fst_node.col_offset), severity='error') return None - if lhs.is_temp and rhs.dtype is NativeVoid(): + if lhs is PyccelSymbol and lhs.is_temp and rhs.dtype is NativeVoid(): return rhs lhs = self._assign_lhs_variable(lhs, d_var, rhs, new_expressions, isinstance(expr, AugAssign), **settings) elif isinstance(lhs, PythonTuple): From 8ef59906f98230ce1b03545e6817e8279ce751d3 Mon Sep 17 00:00:00 2001 From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com> Date: Tue, 10 Jan 2023 19:33:51 +0100 Subject: [PATCH 3/6] Revert cuda_array_managed.py to cuda_main_temp version --- tests/internal/scripts/ccuda/cuda_array_managed.py | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/tests/internal/scripts/ccuda/cuda_array_managed.py b/tests/internal/scripts/ccuda/cuda_array_managed.py index 4e53545929..96898ff506 100644 --- a/tests/internal/scripts/ccuda/cuda_array_managed.py +++ b/tests/internal/scripts/ccuda/cuda_array_managed.py @@ -11,7 +11,7 @@ def square(a): threads_per_block = 5 n_blocks = 1 a = cuda.array([0,1,2,3,4], memory_location = 'managed') - cuda.deviceSynchronize() - square[5, 1](a) - cuda.deviceSynchronize() + cuda.synchronize() + square[n_blocks, threads_per_block](a) + cuda.synchronize() print(a) \ No newline at end of file From 4691db4009395c5563149052c7a724485fcca225 Mon Sep 17 00:00:00 2001 From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com> Date: Tue, 10 Jan 2023 19:46:16 +0100 Subject: [PATCH 4/6] Fix CudaSynchronize printer's name --- 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 e957e3b26e..795822486a 100644 --- a/pyccel/codegen/printing/ccudacode.py +++ b/pyccel/codegen/printing/ccudacode.py @@ -536,7 +536,7 @@ def copy_CudaArray_Data(self, expr): cpy_data = "cudaMemcpy({0}.raw_data, {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): + def _print_CudaSynchronize(self, expr): return 'cudaDeviceSynchronize();\n' def _print_CudaInternalVar(self, expr): From 5f32e4fd0222f299587db3a043bba69379ce0d17 Mon Sep 17 00:00:00 2001 From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com> Date: Tue, 10 Jan 2023 19:52:43 +0100 Subject: [PATCH 5/6] Add test containing only a call to cuda.synchronize --- tests/internal/scripts/ccuda/cuda_synchronize.py | 5 +++++ 1 file changed, 5 insertions(+) create mode 100644 tests/internal/scripts/ccuda/cuda_synchronize.py diff --git a/tests/internal/scripts/ccuda/cuda_synchronize.py b/tests/internal/scripts/ccuda/cuda_synchronize.py new file mode 100644 index 0000000000..c208566846 --- /dev/null +++ b/tests/internal/scripts/ccuda/cuda_synchronize.py @@ -0,0 +1,5 @@ +from pyccel import cuda + +if __name__ == '__main__': + cuda.synchronize() + cuda.synchronize() From 2a851edd0ce5fd41a0251345c43b7ab0b462fd1c Mon Sep 17 00:00:00 2001 From: Aymane Benaissa <47903494+Pinkyboi@users.noreply.github.com> Date: Tue, 10 Jan 2023 21:02:44 +0100 Subject: [PATCH 6/6] Use only one cuda.synchronize in cuda_synchronize.py --- tests/internal/scripts/ccuda/cuda_synchronize.py | 1 - 1 file changed, 1 deletion(-) diff --git a/tests/internal/scripts/ccuda/cuda_synchronize.py b/tests/internal/scripts/ccuda/cuda_synchronize.py index c208566846..bc95e56668 100644 --- a/tests/internal/scripts/ccuda/cuda_synchronize.py +++ b/tests/internal/scripts/ccuda/cuda_synchronize.py @@ -2,4 +2,3 @@ if __name__ == '__main__': cuda.synchronize() - cuda.synchronize()