compiler: restrict nested parallelism to supported compilers (intel)

Author: mloubout

devito/passes/iet/languages/openmp.py

@@ -5,7 +5,9 @@
 from sympy import And, Ne, Not
 
 from devito.arch import AMDGPUX, INTELGPUX, NVIDIAX, PVC
-from devito.arch.compiler import CustomCompiler, GNUCompiler, NvidiaCompiler
+from devito.arch.compiler import (
+    CustomCompiler, GNUCompiler, IntelCompiler, NvidiaCompiler, OneapiCompiler
+)
 from devito.ir import (
     Call, Conditional, DeviceCall, FindSymbols, List, ParallelBlock, PointerCast, Pragma,
     Prodder, While
@@ -276,6 +278,16 @@ def _support_complex_reduction(cls, compiler):
             # Gcc doesn't supports complex reduction
             return not isinstance(compiler, GNUCompiler)
 
+    @classmethod
+    def _support_nested_parallelism(cls, compiler):
+        # In case we have a CustomCompiler
+        if isinstance(compiler, CustomCompiler):
+            compiler = compiler._base()
+        if isinstance(compiler, (IntelCompiler, OneapiCompiler)):  # noqa: SIM103
+            return True
+        else:
+            return False
+
 
 class Ompizer(AbstractOmpizer):
     langbb = OmpBB

devito/passes/iet/parpragma.py

@@ -54,6 +54,10 @@ def _support_array_reduction(cls, compiler):
     def _support_complex_reduction(cls, compiler):
         return False
 
+    @classmethod
+    def _support_nested_parallelism(cls, compiler):
+        return False
+
     @property
     def simd_reg_nbytes(self):
         return self.platform.simd_reg_nbytes
@@ -344,7 +348,8 @@ def _make_guard(self, parregion):
 
     def _make_nested_partree(self, partree):
         # Apply heuristic
-        if self.nhyperthreads <= self.nested:
+        if self.nhyperthreads <= self.nested or \
+           not self._support_nested_parallelism(self.compiler):
             return partree
 
         # Note: there might be multiple sub-trees amenable to nested parallelism,

tests/test_dle.py

@@ -11,6 +11,7 @@
     PrecomputedSparseTimeFunction, ReduceMax, ReduceMin, ReduceMinMax, SpaceDimension,
     SparseTimeFunction, SubDimension, TimeFunction, configuration, cos, dimensions, info
 )
+from devito.arch.compiler import IntelCompiler, OneapiCompiler
 from devito.exceptions import InvalidArgument
 from devito.ir.iet import (
     Expression, FindNodes, IsPerfectIteration, Iteration, retrieve_iteration_tree
@@ -1461,3 +1462,40 @@ def test_collapsing_w_wo_halo(self, exprs, collapsed, scheduling):
 
         assert iterations[1].pragmas[0].ccode.value ==\
             "".join([ompfor_string, scheduling_string])
+
+    @skipif('device')
+    def test_nested_parallelism_support(self):
+        grid = Grid(shape=(10, 10, 10))
+
+        f = Function(name='f', grid=grid, space_order=4)
+        v = TimeFunction(name="v", grid=grid, space_order=4)
+        v1 = TimeFunction(name="v1", grid=grid, space_order=4)
+
+        f.data_with_halo[:] = 0.5
+        v.data_with_halo[:] = 1.
+        v1.data_with_halo[:] = 1.
+
+        eqn = Eq(v.forward, (v.dx * (1 + 2*f) * f).dx)
+        op = Operator(eqn, opt=('advanced', {'openmp': True, 'par-nested': 0}))
+
+        bns, _ = assert_blocking(op, {'x0_blk0'})
+        trees = retrieve_iteration_tree(bns['x0_blk0'])
+        assert len(trees) == 2
+
+        # Check omp pargams
+        assert trees[0][0].pragmas[0].ccode.value == \
+            'omp for collapse(2) schedule(dynamic,1)'
+        if isinstance(configuration['compiler'], (IntelCompiler, OneapiCompiler)):
+            # Supports nested parallelism
+            assert trees[0][2].pragmas[0].ccode.value == \
+                '#pragma omp parallel for collapse(2) schedule(dynamic,1)'\
+                ' num_threads(nthreads_nested)'
+            assert trees[1][2].pragmas[0].ccode.value == \
+                trees[0][2].pragmas[0].ccode.value
+        else:
+            # Most compiler don't support nested parallelism
+            assert not trees[0][2].pragmas
+            assert not trees[1][2].pragmas
+
+        # Should compile properly
+        op.cfunction  # noqa: B018