84 lines · plain
1!===-- module/cooperative_groups.f90 ---------------------------------------===!2!3! Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.4! See https://llvm.org/LICENSE.txt for license information.5! SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception6!7!===------------------------------------------------------------------------===!8 9! CUDA Fortran cooperative groups10 11module cooperative_groups12 13use, intrinsic :: __fortran_builtins, only: c_devptr => __builtin_c_devptr14 15implicit none16 17type :: cluster_group18 type(c_devptr), private :: handle19 integer(4) :: size20 integer(4) :: rank21end type cluster_group22 23type :: grid_group24 type(c_devptr), private :: handle25 integer(4) :: size26 integer(4) :: rank27end type grid_group28 29type :: coalesced_group30 type(c_devptr), private :: handle31 integer(4) :: size32 integer(4) :: rank33end type coalesced_group34 35type :: thread_group36 type(c_devptr), private :: handle37 integer(4) :: size38 integer(4) :: rank39end type thread_group40 41interface42 attributes(device) function cluster_block_index()43 import44 type(dim3) :: cluster_block_index45 end function46end interface47 48interface49 attributes(device) function cluster_dim_blocks()50 import51 type(dim3) :: cluster_dim_blocks52 end function53end interface54 55interface56 attributes(device) function this_cluster()57 import58 type(cluster_group) :: this_cluster59 end function60end interface61 62interface63 attributes(device) function this_grid()64 import65 type(grid_group) :: this_grid66 end function67end interface68 69interface70 attributes(device) function this_thread_block()71 import72 type(thread_group) :: this_thread_block73 end function74end interface75 76interface this_warp77 attributes(device) function this_warp()78 import79 type(coalesced_group) :: this_warp80 end function81end interface82 83end module84