114 lines · plain
1! RUN: bbc -emit-hlfir -fcuda %s -o - | FileCheck %s2! RUN: bbc -emit-hlfir -fcuda %s -o - | fir-opt | FileCheck %s3 4! Test lowering of CUDA kernel loop directive.5 6subroutine sub1()7 integer :: i, j8 integer, parameter :: n = 1009 integer(8) :: istream10 real, device :: a(n), b(n)11 real, device :: c(n,n), d(n,n)12 13! CHECK-LABEL: func.func @_QPsub1()14! CHECK: %[[IV:.*]]:2 = hlfir.declare %{{.*}} {uniq_name = "_QFsub1Ei"} : (!fir.ref<i32>) -> (!fir.ref<i32>, !fir.ref<i32>)15! CHECK: %[[STREAM:.*]]:2 = hlfir.declare %{{.*}} {uniq_name = "_QFsub1Eistream"} : (!fir.ref<i64>) -> (!fir.ref<i64>, !fir.ref<i64>)16! CHECK: %[[IV_J:.*]]:2 = hlfir.declare %{{.*}} {uniq_name = "_QFsub1Ej"} : (!fir.ref<i32>) -> (!fir.ref<i32>, !fir.ref<i32>)17 !$cuf kernel do <<< 1, 2 >>>18 do i = 1, n19 a(i) = a(i) * b(i)20 end do21 22! CHECK: %[[LB:.*]] = fir.convert %c1{{.*}} : (i32) -> index23! CHECK: %[[UB:.*]] = fir.convert %c100{{.*}} : (i32) -> index24! CHECK: %[[STEP:.*]] = arith.constant 1 : index25! CHECK: cuf.kernel<<<%c1_i32, %c2_i32>>> (%[[ARG0:.*]] : index) = (%[[LB]] : index) to (%[[UB]] : index) step (%[[STEP]] : index)26! CHECK-NOT: fir.do_loop27! CHECK: %[[ARG0_I32:.*]] = fir.convert %[[ARG0]] : (index) -> i3228! CHECK: fir.store %[[ARG0_I32]] to %[[IV]]#0 : !fir.ref<i32>29! CHECK: hlfir.assign30 31 32 !$cuf kernel do <<< *, * >>>33 do i = 1, n34 a(i) = a(i) * b(i)35 end do36 37! CHECK: cuf.kernel<<<*, *>>> (%{{.*}} : index) = (%{{.*}} : index) to (%{{.*}} : index) step (%{{.*}} : index)38 39 !$cuf kernel do(2) <<< 1, (256,1) >>>40 do i = 1, n41 do j = 1, n42 c(i,j) = c(i,j) * d(i,j)43 end do44 end do45 46! CHECK: cuf.kernel<<<%c1{{.*}}, (%c256{{.*}}, %c1{{.*}})>>> (%[[ARG0:.*]] : index, %[[ARG1:.*]] : index) = (%{{.*}}, %{{.*}} : index, index) to (%{{.*}}, %{{.*}} : index, index) step (%{{.*}}, %{{.*}} : index, index)47! CHECK: %[[ARG0_I32:.*]] = fir.convert %[[ARG0]] : (index) -> i3248! CHECK: fir.store %[[ARG0_I32]] to %[[IV]]#0 : !fir.ref<i32>49! CHECK: %[[ARG1_I32:.*]] = fir.convert %[[ARG1]] : (index) -> i3250! CHECK: fir.store %[[ARG1_I32]] to %[[IV_J]]#0 : !fir.ref<i32>51! CHECK: {n = 2 : i64}52 53 !$cuf kernel do(2) <<< (1,*), (256,1) >>>54 do i = 1, n55 do j = 1, n56 c(i,j) = c(i,j) * d(i,j)57 end do58 end do59! CHECK: cuf.kernel<<<(%c1{{.*}}, %c0{{.*}}), (%c256{{.*}}, %c1{{.*}})>>> (%{{.*}} : index, %{{.*}} : index) = (%{{.*}}, %{{.*}} : index, index) to (%{{.*}}, %{{.*}} : index, index) step (%{{.*}}, %{{.*}} : index, index)60 61!$cuf kernel do(2) <<< (*,*), (32,4) >>>62 do i = 1, n63 do j = 1, n64 c(i,j) = c(i,j) * d(i,j)65 end do66 end do67 68! CHECK: cuf.kernel<<<*, (%c32{{.*}}, %c4{{.*}})>>> (%{{.*}} : index, %{{.*}} : index) = (%{{.*}}, %{{.*}} : index, index) to (%{{.*}}, %{{.*}} : index, index) step (%{{.*}}, %{{.*}} : index, index)69 70 !$cuf kernel do(2) <<< (*,*), (*,*), stream=istream >>>71 do i = 1, n72 do j = 1, n73 c(i,j) = c(i,j) * d(i,j)74 end do75 end do76end77 78! CHECK: cuf.kernel<<<*, *, stream = %[[STREAM]]#0 : !fir.ref<i64>>>>79 80 81! Test lowering with unstructured construct inside.82subroutine sub2(m,a,b)83 integer :: m84 real, device :: a(m,m), b(m)85 integer :: i,j86 !$cuf kernel do<<<*,*>>>87 88 do j = 1, m89 i = 190 do while (a(i,j).eq.0)91 i = i + 192 end do93 b(j) = i94 end do95end subroutine96 97! CHECK-LABEL: func.func @_QPsub298! CHECK: cuf.kernel99 100subroutine sub3()101 integer, device :: a(10), b(10)102 integer :: lb = 1103 integer :: n = 10104 integer :: s = 1105 106 !$cuf kernel do <<< *, * >>>107 do i = lb, n, s108 a(i) = a(i) * b(i)109 end do110end111 112! CHECK-LABEL: func.func @_QPsub3113! CHECK: cuf.kernel114