56 lines · plain
1! RUN: bbc -emit-hlfir -fcuda %s -o - | FileCheck %s2 3attributes(global) subroutine test_this_cluster()4 use cooperative_groups5 type(cluster_group) :: cluster6 7 cluster = this_cluster()8end subroutine9 10! CHECK-LABEL: func.func @_QPtest_this_cluster() attributes {cuf.proc_attr = #cuf.cuda_proc<global>}11! CHECK: %{{.*}} = fir.alloca !fir.type<_QMcooperative_groupsTcluster_group12! CHECK: %[[RES:.*]] = fir.alloca !fir.type<_QMcooperative_groupsTcluster_group{_QMcooperative_groupsTcluster_group.handle:!fir.type<_QM__fortran_builtinsT__builtin_c_devptr{cptr:!fir.type<_QM__fortran_builtinsT__builtin_c_ptr{__address:i64}>}>,size:i32,rank:i32}>13! CHECK: %[[RANK:.*]] = nvvm.read.ptx.sreg.cluster.ctarank : i3214! CHECK: %[[RANK_1:.*]] = arith.addi %[[RANK]], %c1{{.*}} : i32 15! CHECK: %[[RANK_COORD:.*]] = fir.coordinate_of %[[RES]], rank : (!fir.ref<!fir.type<_QMcooperative_groupsTcluster_group{_QMcooperative_groupsTcluster_group.handle:!fir.type<_QM__fortran_builtinsT__builtin_c_devptr{cptr:!fir.type<_QM__fortran_builtinsT__builtin_c_ptr{__address:i64}>}>,size:i32,rank:i32}>>) -> !fir.ref<i32>16! CHECK: fir.store %[[RANK_1]] to %[[RANK_COORD]] : !fir.ref<i32>17 18attributes(global) subroutine test_cluster_dim_blocks()19 use cooperative_groups20 type(dim3) :: clusterDim21 22 clusterDim = cluster_dim_blocks()23end subroutine24 25! CHECK-LABEL: func.func @_QPtest_cluster_dim_blocks() attributes {cuf.proc_attr = #cuf.cuda_proc<global>}26! CHECK: %[[X:.*]] = nvvm.read.ptx.sreg.cluster.nctaid.x : i3227! CHECK: %[[COORD_X:.*]] = fir.coordinate_of %{{.*}}, x : (!fir.ref<!fir.type<_QM__fortran_builtinsT__builtin_dim3{x:i32,y:i32,z:i32}>>) -> !fir.ref<i32>28! CHECK: fir.store %[[X]] to %[[COORD_X]] : !fir.ref<i32>29! CHECK: %[[Y:.*]] = nvvm.read.ptx.sreg.cluster.nctaid.y : i3230! CHECK: %[[COORD_Y:.*]] = fir.coordinate_of %{{.*}}, y : (!fir.ref<!fir.type<_QM__fortran_builtinsT__builtin_dim3{x:i32,y:i32,z:i32}>>) -> !fir.ref<i32>31! CHECK: fir.store %[[Y]] to %[[COORD_Y]] : !fir.ref<i32>32! CHECK: %[[Z:.*]] = nvvm.read.ptx.sreg.cluster.nctaid.z : i3233! CHECK: %[[COORD_Z:.*]] = fir.coordinate_of %{{.*}}, z : (!fir.ref<!fir.type<_QM__fortran_builtinsT__builtin_dim3{x:i32,y:i32,z:i32}>>) -> !fir.ref<i32>34! CHECK: fir.store %[[Z]] to %[[COORD_Z]] : !fir.ref<i32>35 36attributes(global) subroutine test_cluster_block_index()37 use cooperative_groups38 type(dim3) :: blockIndex39 40 blockIndex = cluster_block_index()41end subroutine42 43! CHECK-LABEL: func.func @_QPtest_cluster_block_index() attributes {cuf.proc_attr = #cuf.cuda_proc<global>}44! CHECK: %[[X:.*]] = nvvm.read.ptx.sreg.cluster.ctaid.x : i3245! CHECK: %[[X1:.*]] = arith.addi %[[X]], %c1{{.*}} : i32 46! CHECK: %[[COORD_X:.*]] = fir.coordinate_of %{{.*}}, x : (!fir.ref<!fir.type<_QM__fortran_builtinsT__builtin_dim3{x:i32,y:i32,z:i32}>>) -> !fir.ref<i32>47! CHECK: fir.store %[[X1]] to %[[COORD_X]] : !fir.ref<i32>48! CHECK: %[[Y:.*]] = nvvm.read.ptx.sreg.cluster.ctaid.y : i3249! CHECK: %[[Y1:.*]] = arith.addi %[[Y]], %c1{{.*}} : i32 50! CHECK: %[[COORD_Y:.*]] = fir.coordinate_of %{{.*}}, y : (!fir.ref<!fir.type<_QM__fortran_builtinsT__builtin_dim3{x:i32,y:i32,z:i32}>>) -> !fir.ref<i32>51! CHECK: fir.store %[[Y1]] to %[[COORD_Y]] : !fir.ref<i32>52! CHECK: %[[Z:.*]] = nvvm.read.ptx.sreg.cluster.ctaid.z : i3253! CHECK: %[[Z1:.*]] = arith.addi %[[Z]], %c1{{.*}} : i32 54! CHECK: %[[COORD_Z:.*]] = fir.coordinate_of %{{.*}}, z : (!fir.ref<!fir.type<_QM__fortran_builtinsT__builtin_dim3{x:i32,y:i32,z:i32}>>) -> !fir.ref<i32>55! CHECK: fir.store %[[Z1]] to %[[COORD_Z]] : !fir.ref<i32>56