[flang-commits] [flang] Allow do concurrent inside cuf kernel directive (PR #127693)

[Zhen Wang via flang-commits]

================
@@ -0,0 +1,39 @@
+! RUN: bbc -emit-hlfir -fcuda %s -o - | FileCheck %s
+
+! Check if do concurrent works inside cuf kernel directive
+
+subroutine doconc1
+  integer :: i, n
+  integer, managed :: a(3)
+  a(:) = -1
+  n = 3
+  n = n - 1
+  !$cuf kernel do
+  do concurrent(i=1:n)
+    a(i) = 1
+  end do
+end
+
+! CHECK: func.func @_QPdoconc1() {
+! CHECK: %[[DECL:.*]]:2 = hlfir.declare %{{.*}}#0 {uniq_name = "_QFdoconc1Ei"} : (!fir.ref<i32>) -> (!fir.ref<i32>, !fir.ref<i32>)
+! CHECK: cuf.kernel<<<*, *>>>
+! CHECK: %{{.*}} = fir.load %[[DECL]]#0 : !fir.ref<i32>
+
+subroutine doconc2
+  integer :: i, j, m, n
+  integer, managed :: a(2, 4)
+  m = 2
+  n = 4
+  a(:,:) = -1
+  !$cuf kernel do
+  do concurrent(i=1:m,j=1:n)
+    a(i,j) = i+j
+  end do
+end
+
+! CHECK: func.func @_QPdoconc2() {
+! CHECK: %[[DECLI:.*]]:2 = hlfir.declare %{{.*}}#0 {uniq_name = "_QFdoconc2Ei"} : (!fir.ref<i32>) -> (!fir.ref<i32>, !fir.ref<i32>)
+! CHECK: %[[DECLJ:.*]]:2 = hlfir.declare %{{.*}}#0 {uniq_name = "_QFdoconc2Ej"} : (!fir.ref<i32>) -> (!fir.ref<i32>, !fir.ref<i32>)
+! CHECK: cuf.kernel<<<*, *>>> (%arg0 : i32, %arg1 : i32) = (%{{.*}}, %{{.*}} : index, index) to (%{{.*}}, %{{.*}} : index, index)  step (%{{.*}}, %{{.*}} : index, index) {
----------------
wangzpgi wrote:

> Is this really the semantic of `!$cuf kernel do`? For me it means `!$cuf kernel do(1)` so only the first range is part of the cuf kernel operation and the rest should be nested inside. Let me know if the semantic is different.

Yes, it's the same semantic with `do` and `do(1)`, we have this multi range with `cuf.kernel`.

https://github.com/llvm/llvm-project/pull/127693