753 lines · plain
1! RUN: bbc -emit-hlfir -fcuda %s -o - | FileCheck %s2 3! Test CUDA Fortran procedures available in cudadevice module4 5attributes(global) subroutine devsub()6 implicit none7 integer :: ret8 real(2) :: r29 real(4) :: af10 real(8) :: ad11 integer(4) :: ai12 integer(8) :: al13 integer(8) :: time14 integer :: smalltime15 integer(4) :: res, offset16 integer(8) :: resl17 real(2) :: r2a(2)18 real(2) :: tmp2(2)19 20 integer :: tid21 tid = threadIdx%x22 23 call syncthreads()24 call syncwarp(1)25 ret = syncthreads_and(1)26 res = syncthreads_and(tid > offset)27 ret = syncthreads_count(1)28 ret = syncthreads_count(tid > offset)29 ret = syncthreads_or(1)30 ret = syncthreads_or(tid > offset)31 32 ai = atomicadd(ai, 1_4)33 al = atomicadd(al, 1_8)34 af = atomicadd(af, 1.0_4)35 ad = atomicadd(ad, 1.0_8)36 ai = atomicadd(r2a, tmp2)37 38 ai = atomicsub(ai, 1_4)39 al = atomicsub(al, 1_8)40 af = atomicsub(af, 1.0_4)41 ad = atomicsub(ad, 1.0_8)42 43 ai = atomicmax(ai, 1_4)44 al = atomicmax(al, 1_8)45 af = atomicmax(af, 1.0_4)46 ad = atomicmax(ad, 1.0_8)47 48 ai = atomicmin(ai, 1_4)49 al = atomicmin(al, 1_8)50 af = atomicmin(af, 1.0_4)51 ad = atomicmin(ad, 1.0_8)52 53 ai = atomicand(ai, 1_4)54 ai = atomicor(ai, 1_4)55 ai = atomicinc(ai, 1_4)56 ai = atomicdec(ai, 1_4)57 58 smalltime = clock()59 time = clock64()60 time = globalTimer()61 62 res = __popc(ai)63 res = __popc(al)64 res = __ffs(ai)65 res = __ffs(al)66 res = __brev(ai)67 resl = __brev(al)68 69 res = __clz(ai)70 res = __clz(al) 71 af = __cosf(af)72 ad = __ddiv_rn(ad, ad)73 ad = __ddiv_rz(ad, ad)74 ad = __ddiv_ru(ad, ad)75 ad = __ddiv_rd(ad, ad)76 af = __double2float_rn(ad)77 af = __double2float_rz(ad)78 af = __double2float_ru(ad)79 af = __double2float_rd(ad)80 ai = __double2int_rd(ad)81 ai = __double2int_rn(ad)82 ai = __double2int_ru(ad)83 ai = __double2int_rz(ad)84 ai = __double2uint_rd(ad)85 ai = __double2uint_rn(ad)86 ai = __double2uint_ru(ad)87 ai = __double2uint_rz(ad)88 ai = __mul24(ai, ai)89 ai = __umul24(ai, ai)90 af = __powf(af, af)91 ad = __ull2double_rd(al)92 ad = __ull2double_rn(al)93 ad = __ull2double_ru(al)94 ad = __ull2double_rz(al)95 r2 = __float2half_rn(af)96 af = __half2float(r2)97 ad = __ll2double_rd(al)98 ad = __ll2double_rn(al)99 ad = __ll2double_ru(al)100 ad = __ll2double_rz(al)101end102 103! CHECK-LABEL: func.func @_QPdevsub() attributes {cuf.proc_attr = #cuf.cuda_proc<global>}104! CHECK: nvvm.barrier0105! CHECK: nvvm.bar.warp.sync %c1{{.*}} : i32 106! CHECK: %{{.*}} = nvvm.barrier #nvvm.reduction<and> %c1{{.*}} -> i32107! CHECK: %[[A:.*]] = fir.load %{{.*}} : !fir.ref<i32>108! CHECK: %[[B:.*]] = fir.load %{{.*}} : !fir.ref<i32>109! CHECK: %[[CMP:.*]] = arith.cmpi sgt, %[[A]], %[[B]] : i32110! CHECK: %[[CONV:.*]] = fir.convert %[[CMP]] : (i1) -> i32111! CHECK: %{{.*}} = nvvm.barrier #nvvm.reduction<and> %[[CONV]] -> i32112! CHECK: %{{.*}} = nvvm.barrier #nvvm.reduction<popc> %c1{{.*}} -> i32113! CHECK: %[[A:.*]] = fir.load %{{.*}} : !fir.ref<i32>114! CHECK: %[[B:.*]] = fir.load %{{.*}} : !fir.ref<i32>115! CHECK: %[[CMP:.*]] = arith.cmpi sgt, %[[A]], %[[B]] : i32116! CHECK: %[[CONV:.*]] = fir.convert %[[CMP]] : (i1) -> i32117! CHECK: %{{.*}} = nvvm.barrier #nvvm.reduction<popc> %[[CONV]] -> i32118! CHECK: %{{.*}} = nvvm.barrier #nvvm.reduction<or> %c1{{.*}} -> i32119! CHECK: %[[A:.*]] = fir.load %{{.*}} : !fir.ref<i32>120! CHECK: %[[B:.*]] = fir.load %{{.*}} : !fir.ref<i32>121! CHECK: %[[CMP:.*]] = arith.cmpi sgt, %[[A]], %[[B]] : i32122! CHECK: %[[CONV:.*]] = fir.convert %[[CMP]] : (i1) -> i32123! CHECK: %{{.*}} = nvvm.barrier #nvvm.reduction<or> %[[CONV]] -> i32124! CHECK: %{{.*}} = llvm.atomicrmw add %{{.*}}, %{{.*}} seq_cst : !llvm.ptr, i32125! CHECK: %{{.*}} = llvm.atomicrmw add %{{.*}}, %{{.*}} seq_cst : !llvm.ptr, i64126! CHECK: %{{.*}} = llvm.atomicrmw fadd %{{.*}}, %{{.*}} seq_cst : !llvm.ptr, f32127! CHECK: %{{.*}} = llvm.atomicrmw fadd %{{.*}}, %{{.*}} seq_cst : !llvm.ptr, f64128! CHECK: %{{.*}} = llvm.atomicrmw fadd %{{.*}}, %{{.*}} seq_cst : !llvm.ptr, vector<2xf16>129 130! CHECK: %{{.*}} = llvm.atomicrmw sub %{{.*}}, %{{.*}} seq_cst : !llvm.ptr, i32131! CHECK: %{{.*}} = llvm.atomicrmw sub %{{.*}}, %{{.*}} seq_cst : !llvm.ptr, i64132! CHECK: %{{.*}} = llvm.atomicrmw fsub %{{.*}}, %{{.*}} seq_cst : !llvm.ptr, f32133! CHECK: %{{.*}} = llvm.atomicrmw fsub %{{.*}}, %{{.*}} seq_cst : !llvm.ptr, f64134 135! CHECK: %{{.*}} = llvm.atomicrmw max %{{.*}}, %{{.*}} seq_cst : !llvm.ptr, i32136! CHECK: %{{.*}} = llvm.atomicrmw max %{{.*}}, %{{.*}} seq_cst : !llvm.ptr, i64137! CHECK: %{{.*}} = llvm.atomicrmw fmax %{{.*}}, %{{.*}} seq_cst : !llvm.ptr, f32138! CHECK: %{{.*}} = llvm.atomicrmw fmax %{{.*}}, %{{.*}} seq_cst : !llvm.ptr, f64139 140! CHECK: %{{.*}} = llvm.atomicrmw min %{{.*}}, %{{.*}} seq_cst : !llvm.ptr, i32141! CHECK: %{{.*}} = llvm.atomicrmw min %{{.*}}, %{{.*}} seq_cst : !llvm.ptr, i64142! CHECK: %{{.*}} = llvm.atomicrmw fmin %{{.*}}, %{{.*}} seq_cst : !llvm.ptr, f32143! CHECK: %{{.*}} = llvm.atomicrmw fmin %{{.*}}, %{{.*}} seq_cst : !llvm.ptr, f64144 145! CHECK: %{{.*}} = llvm.atomicrmw _and %{{.*}}, %{{.*}} seq_cst : !llvm.ptr, i32146! CHECK: %{{.*}} = llvm.atomicrmw _or %{{.*}}, %{{.*}} seq_cst : !llvm.ptr, i32147! CHECK: %{{.*}} = llvm.atomicrmw uinc_wrap %{{.*}}, %{{.*}} seq_cst : !llvm.ptr, i32148! CHECK: %{{.*}} = llvm.atomicrmw udec_wrap %{{.*}}, %{{.*}} seq_cst : !llvm.ptr, i32149 150! CHECK: %{{.*}} = nvvm.read.ptx.sreg.clock : i32151! CHECK: %{{.*}} = nvvm.read.ptx.sreg.clock64 : i64152! CHECK: %{{.*}} = nvvm.read.ptx.sreg.globaltimer : i64153 154! CHECK: %{{.*}} = fir.call @__nv_popc(%{{.*}}) proc_attrs<bind_c> fastmath<contract> : (i32) -> i32155! CHECK: %{{.*}} = fir.call @__nv_popcll(%{{.*}}) proc_attrs<bind_c> fastmath<contract> : (i64) -> i32156! CHECK: %{{.*}} = fir.call @__nv_ffs(%{{.*}}) proc_attrs<bind_c> fastmath<contract> : (i32) -> i32157! CHECK: %{{.*}} = fir.call @__nv_ffsll(%{{.*}}) proc_attrs<bind_c> fastmath<contract> : (i64) -> i32158! CHECK: %{{.*}} = fir.call @__nv_brev(%{{.*}}) proc_attrs<bind_c> fastmath<contract> : (i32) -> i32159! CHECK: %{{.*}} = fir.call @__nv_brevll(%{{.*}}) proc_attrs<bind_c> fastmath<contract> : (i64) -> i64160! CHECK: %{{.*}} = fir.call @__nv_clz(%{{.*}}) proc_attrs<bind_c> fastmath<contract> : (i32) -> i32161! CHECK: %{{.*}} = fir.call @__nv_clzll(%{{.*}}) proc_attrs<bind_c> fastmath<contract> : (i64) -> i32162! CHECK: %{{.*}} = fir.call @__nv_fast_cosf(%{{.*}}) proc_attrs<bind_c> fastmath<contract> : (f32) -> f32163! CHECK: %{{.*}} = fir.call @__nv_ddiv_rn(%{{.*}}, %{{.*}}) proc_attrs<bind_c> fastmath<contract> : (f64, f64) -> f64164! CHECK: %{{.*}} = fir.call @__nv_ddiv_rz(%{{.*}}, %{{.*}}) proc_attrs<bind_c> fastmath<contract> : (f64, f64) -> f64165! CHECK: %{{.*}} = fir.call @__nv_ddiv_ru(%{{.*}}, %{{.*}}) proc_attrs<bind_c> fastmath<contract> : (f64, f64) -> f64166! CHECK: %{{.*}} = fir.call @__nv_ddiv_rd(%{{.*}}, %{{.*}}) proc_attrs<bind_c> fastmath<contract> : (f64, f64) -> f64167! CHECK: %{{.*}} = fir.call @__nv_double2float_rn(%{{.*}}) proc_attrs<bind_c> fastmath<contract> : (f64) -> f32168! CHECK: %{{.*}} = fir.call @__nv_double2float_rz(%{{.*}}) proc_attrs<bind_c> fastmath<contract> : (f64) -> f32169! CHECK: %{{.*}} = fir.call @__nv_double2float_ru(%{{.*}}) proc_attrs<bind_c> fastmath<contract> : (f64) -> f32170! CHECK: %{{.*}} = fir.call @__nv_double2float_rd(%{{.*}}) proc_attrs<bind_c> fastmath<contract> : (f64) -> f32171! CHECK: %{{.*}} = fir.call @__nv_double2int_rd(%{{.*}}) proc_attrs<bind_c> fastmath<contract> : (f64) -> i32172! CHECK: %{{.*}} = fir.call @__nv_double2int_rn(%{{.*}}) proc_attrs<bind_c> fastmath<contract> : (f64) -> i32173! CHECK: %{{.*}} = fir.call @__nv_double2int_ru(%{{.*}}) proc_attrs<bind_c> fastmath<contract> : (f64) -> i32174! CHECK: %{{.*}} = fir.call @__nv_double2int_rz(%{{.*}}) proc_attrs<bind_c> fastmath<contract> : (f64) -> i32175! CHECK: %{{.*}} = fir.call @__nv_double2uint_rd(%{{.*}}) proc_attrs<bind_c> fastmath<contract> : (f64) -> i32176! CHECK: %{{.*}} = fir.call @__nv_double2uint_rn(%{{.*}}) proc_attrs<bind_c> fastmath<contract> : (f64) -> i32177! CHECK: %{{.*}} = fir.call @__nv_double2uint_ru(%{{.*}}) proc_attrs<bind_c> fastmath<contract> : (f64) -> i32178! CHECK: %{{.*}} = fir.call @__nv_double2uint_rz(%{{.*}}) proc_attrs<bind_c> fastmath<contract> : (f64) -> i32179! CHECK: %{{.*}} = fir.call @__nv_mul24(%{{.*}}, %{{.*}}) proc_attrs<bind_c> fastmath<contract> : (i32, i32) -> i32180! CHECK: %{{.*}} = fir.call @__nv_umul24(%{{.*}}, %{{.*}}) proc_attrs<bind_c> fastmath<contract> : (i32, i32) -> i32181! CHECK: %{{.*}} = fir.call @__nv_fast_powf(%{{.*}}, %{{.*}}) proc_attrs<bind_c> fastmath<contract> : (f32, f32) -> f32182! CHECK: %{{.*}} = fir.call @__nv_ull2double_rd(%{{.*}}) proc_attrs<bind_c> fastmath<contract> : (i64) -> f64183! CHECK: %{{.*}} = fir.call @__nv_ull2double_rn(%{{.*}}) proc_attrs<bind_c> fastmath<contract> : (i64) -> f64184! CHECK: %{{.*}} = fir.call @__nv_ull2double_ru(%{{.*}}) proc_attrs<bind_c> fastmath<contract> : (i64) -> f64185! CHECK: %{{.*}} = fir.call @__nv_ull2double_rz(%{{.*}}) proc_attrs<bind_c> fastmath<contract> : (i64) -> f64186! CHECK: %{{.*}} = fir.call @__nv_float2half_rn(%{{.*}}) proc_attrs<bind_c> fastmath<contract> : (f32) -> f16187! CHECK: %{{.*}} = fir.call @__nv_half2float(%{{.*}}) proc_attrs<bind_c> fastmath<contract> : (f16) -> f32188! CHECK: %{{.*}} = fir.call @__nv_ll2double_rd(%{{.*}}) proc_attrs<bind_c> fastmath<contract> : (i64) -> f64189! CHECK: %{{.*}} = fir.call @__nv_ll2double_rn(%{{.*}}) proc_attrs<bind_c> fastmath<contract> : (i64) -> f64190! CHECK: %{{.*}} = fir.call @__nv_ll2double_ru(%{{.*}}) proc_attrs<bind_c> fastmath<contract> : (i64) -> f64191! CHECK: %{{.*}} = fir.call @__nv_ll2double_rz(%{{.*}}) proc_attrs<bind_c> fastmath<contract> : (i64) -> f64192 193subroutine host1()194 integer, device :: a(32)195 integer, device :: ret196 integer :: i, j197 198block; use cudadevice199 !$cuf kernel do(1) <<<*,32>>>200 do i = 1, 32201 a(i) = a(i) * 2.0202 call syncthreads()203 a(i) = a(i) + a(j) - 34.0204 205 call syncwarp(1)206 ret = syncthreads_and(1)207 ret = syncthreads_count(1)208 ret = syncthreads_or(1)209 end do210end block211end 212 213! CHECK-LABEL: func.func @_QPhost1()214! CHECK: cuf.kernel215! CHECK: nvvm.barrier0216! CHECK: nvvm.bar.warp.sync %c1{{.*}} : i32 217! CHECK: nvvm.barrier #nvvm.reduction<and> %c1{{.*}} -> i32218! CHECK: nvvm.barrier #nvvm.reduction<popc> %c1{{.*}} -> i32219! CHECK: nvvm.barrier #nvvm.reduction<or> %c1{{.*}} -> i32220 221attributes(device) subroutine testMatch()222 integer :: a, ipred, mask, v32223 integer(8) :: v64224 real(4) :: r4225 real(8) :: r8226 a = match_all_sync(mask, v32, ipred)227 a = match_all_sync(mask, v64, ipred)228 a = match_all_sync(mask, r4, ipred)229 a = match_all_sync(mask, r8, ipred)230end subroutine231 232! CHECK-LABEL: func.func @_QPtestmatch()233! CHECK: %{{.*}} = nvvm.match.sync all %{{.*}}, %{{.*}} : i32 -> !llvm.struct<(i32, i1)>234! CHECK: %{{.*}} = nvvm.match.sync all %{{.*}}, %{{.*}} : i64 -> !llvm.struct<(i32, i1)>235! CHECK: %{{.*}} = nvvm.match.sync all %{{.*}}, %{{.*}} : i32 -> !llvm.struct<(i32, i1)>236! CHECK: %{{.*}} = nvvm.match.sync all %{{.*}}, %{{.*}} : i64 -> !llvm.struct<(i32, i1)>237 238attributes(device) subroutine testMatchAny()239 integer :: a, mask, v32240 integer(8) :: v64241 real(4) :: r4242 real(8) :: r8243 a = match_any_sync(mask, v32)244 a = match_any_sync(mask, v64)245 a = match_any_sync(mask, r4)246 a = match_any_sync(mask, r8)247end subroutine248 249! CHECK-LABEL: func.func @_QPtestmatchany()250! CHECK: %{{.*}} = nvvm.match.sync any %{{.*}}, %{{.*}} : i32 -> i32251! CHECK: %{{.*}} = nvvm.match.sync any %{{.*}}, %{{.*}} : i64 -> i32252! CHECK: %{{.*}} = nvvm.match.sync any %{{.*}}, %{{.*}} : i32 -> i32 253! CHECK: %{{.*}} = nvvm.match.sync any %{{.*}}, %{{.*}} : i64 -> i32254 255attributes(device) subroutine testAtomic(aa, n)256 integer :: aa(*)257 integer, intent(in) :: n258 integer :: a, istat, j, i259 real :: r260 istat = atomicexch(a,0)261 istat = atomicexch(r, 0.0)262 istat = atomicxor(a, j)263 istat = atomiccas(a, i, 14)264 do i = 1, n265 istat = atomicxor(aa, i)266 istat = atomiccas(aa, i, 14)267 istat = atomicexch(aa, 0)268 end do269end subroutine270 271! CHECK-LABEL: func.func @_QPtestatomic272! CHECK: llvm.atomicrmw xchg %{{.*}}, %c0{{.*}} seq_cst : !llvm.ptr, i32273! CHECK: llvm.atomicrmw xchg %{{.*}}, %{{.*}} seq_cst : !llvm.ptr, f32274! CHECK: llvm.atomicrmw _xor %{{.*}}, %{{.*}} seq_cst : !llvm.ptr, i32275! CHECK: %[[ADDR:.*]] = builtin.unrealized_conversion_cast %{{.*}}#0 : !fir.ref<i32> to !llvm.ptr276! CHECK: llvm.cmpxchg %[[ADDR]], %{{.*}}, %c14{{.*}} acq_rel monotonic : !llvm.ptr, i32277! CHECK: fir.do_loop278! CHECK: llvm.atomicrmw _xor %{{.*}}, %{{.*}} seq_cst : !llvm.ptr, i32279! CHECK: %[[ADDR:.*]] = builtin.unrealized_conversion_cast %{{.*}}#1 : !fir.ref<!fir.array<?xi32>> to !llvm.ptr 280! CHECK: llvm.cmpxchg %[[ADDR]], %{{.*}}, %c14{{.*}} acq_rel monotonic : !llvm.ptr, i32281! CHECK: llvm.atomicrmw xchg %{{.*}}, %c0{{.*}} seq_cst : !llvm.ptr, i32282 283attributes(device) subroutine testAtomic2()284 integer(8) :: a, i, istat285 istat = atomiccas(a, i, 14)286end subroutine287 288! CHECK-LABEL: func.func @_QPtestatomic2()289! CHECK: %[[VAL:.*]] = fir.convert %c14{{.*}} : (i32) -> i64290! CHECK: %[[ADDR:.*]] = builtin.unrealized_conversion_cast %{{.*}}#0 : !fir.ref<i64> to !llvm.ptr291! CHECK: llvm.cmpxchg %{{.*}}, %{{.*}}, %[[VAL]] acq_rel monotonic : !llvm.ptr, i64292 293attributes(device) subroutine testAtomic3()294 real :: a, i, istat295 istat = atomiccas(a, i, 14.0)296end subroutine297 298! CHECK-LABEL: func.func @_QPtestatomic3()299! CHECK: %[[BCAST1:.*]] = llvm.bitcast %{{.*}} : f32 to i32300! CHECK: %[[BCAST2:.*]] = llvm.bitcast %{{.*}} : f32 to i32301! CHECK: %[[CAST:.*]] = builtin.unrealized_conversion_cast %{{.*}}#0 : !fir.ref<f32> to !llvm.ptr302! CHECK: llvm.cmpxchg %[[CAST]], %[[BCAST1]], %[[BCAST2]] acq_rel monotonic : !llvm.ptr, i32303 304attributes(device) subroutine testAtomic4()305 real(8) :: a, i, istat306 istat = atomiccas(a, i, 14.0d0)307end subroutine308 309! CHECK-LABEL: func.func @_QPtestatomic4()310! CHECK: %[[BCAST1:.*]] = llvm.bitcast %{{.*}} : f64 to i64311! CHECK: %[[BCAST2:.*]] = llvm.bitcast %{{.*}} : f64 to i64312! CHECK: %[[CAST:.*]] = builtin.unrealized_conversion_cast %{{.*}}#0 : !fir.ref<f64> to !llvm.ptr313! CHECK: %[[ATOMIC:.*]] = llvm.cmpxchg %[[CAST]], %[[BCAST1]], %[[BCAST2]] acq_rel monotonic : !llvm.ptr, i64314! CHECK: %[[RES:.*]] = llvm.extractvalue %[[ATOMIC]][1] : !llvm.struct<(i64, i1)> 315 316attributes(global) subroutine __ldXXi4(b)317 integer, device :: b(*)318 integer, device :: x(4)319 x(1:4) = __ldca(b(i:j))320 x = __ldcg(b(i:j))321 x = __ldcs(b(i:j))322 x(1:4) = __ldlu(b(i:j))323 x(1:4) = __ldcv(b(i:j))324end325 326! CHECK-LABEL: func.func @_QP__ldxxi4327! CHECK: fir.call @__ldca_i4x4_(%{{.*}}, %{{.*}}) fastmath<contract> : (!fir.ref<!fir.array<4xi32>>, !fir.ref<!fir.array<?xi32>>) -> ()328! CHECK: fir.call @__ldcg_i4x4_(%{{.*}}, %{{.*}}) fastmath<contract> : (!fir.ref<!fir.array<4xi32>>, !fir.ref<!fir.array<?xi32>>) -> ()329! CHECK: fir.call @__ldcs_i4x4_(%{{.*}}, %{{.*}}) fastmath<contract> : (!fir.ref<!fir.array<4xi32>>, !fir.ref<!fir.array<?xi32>>) -> ()330! CHECK: fir.call @__ldlu_i4x4_(%{{.*}}, %{{.*}}) fastmath<contract> : (!fir.ref<!fir.array<4xi32>>, !fir.ref<!fir.array<?xi32>>) -> ()331! CHECK: fir.call @__ldcv_i4x4_(%{{.*}}, %{{.*}}) fastmath<contract> : (!fir.ref<!fir.array<4xi32>>, !fir.ref<!fir.array<?xi32>>) -> ()332 333attributes(global) subroutine __ldXXi8(b)334 integer(8), device :: b(*)335 integer(8), device :: x(2)336 x(1:2) = __ldca(b(i:j))337 x = __ldcg(b(i:j))338 x = __ldcs(b(i:j))339 x(1:2) = __ldlu(b(i:j))340 x(1:2) = __ldcv(b(i:j))341end342 343! CHECK-LABEL: func.func @_QP__ldxxi8344! CHECK: fir.call @__ldca_i8x2_(%{{.*}}, %{{.*}}) fastmath<contract> : (!fir.ref<!fir.array<2xi64>>, !fir.ref<!fir.array<?xi64>>) -> ()345! CHECK: fir.call @__ldcg_i8x2_(%{{.*}}, %{{.*}}) fastmath<contract> : (!fir.ref<!fir.array<2xi64>>, !fir.ref<!fir.array<?xi64>>) -> ()346! CHECK: fir.call @__ldcs_i8x2_(%{{.*}}, %{{.*}}) fastmath<contract> : (!fir.ref<!fir.array<2xi64>>, !fir.ref<!fir.array<?xi64>>) -> ()347! CHECK: fir.call @__ldlu_i8x2_(%{{.*}}, %{{.*}}) fastmath<contract> : (!fir.ref<!fir.array<2xi64>>, !fir.ref<!fir.array<?xi64>>) -> ()348! CHECK: fir.call @__ldcv_i8x2_(%{{.*}}, %{{.*}}) fastmath<contract> : (!fir.ref<!fir.array<2xi64>>, !fir.ref<!fir.array<?xi64>>) -> ()349 350attributes(global) subroutine __ldXXr4(b)351 real, device :: b(*)352 real, device :: x(4)353 x(1:4) = __ldca(b(i:j))354 x = __ldcg(b(i:j))355 x = __ldcs(b(i:j))356 x(1:4) = __ldlu(b(i:j))357 x(1:4) = __ldcv(b(i:j))358end359 360! CHECK-LABEL: func.func @_QP__ldxxr4361! CHECK: fir.call @__ldca_r4x4_(%{{.*}}, %{{.*}}) fastmath<contract> : (!fir.ref<!fir.array<4xf32>>, !fir.ref<!fir.array<?xf32>>) -> ()362! CHECK: fir.call @__ldcg_r4x4_(%{{.*}}, %{{.*}}) fastmath<contract> : (!fir.ref<!fir.array<4xf32>>, !fir.ref<!fir.array<?xf32>>) -> ()363! CHECK: fir.call @__ldcs_r4x4_(%{{.*}}, %{{.*}}) fastmath<contract> : (!fir.ref<!fir.array<4xf32>>, !fir.ref<!fir.array<?xf32>>) -> ()364! CHECK: fir.call @__ldlu_r4x4_(%{{.*}}, %{{.*}}) fastmath<contract> : (!fir.ref<!fir.array<4xf32>>, !fir.ref<!fir.array<?xf32>>) -> ()365! CHECK: fir.call @__ldcv_r4x4_(%{{.*}}, %{{.*}}) fastmath<contract> : (!fir.ref<!fir.array<4xf32>>, !fir.ref<!fir.array<?xf32>>) -> ()366 367attributes(global) subroutine __ldXXr2(b)368 real(2), device :: b(*)369 real(2), device :: x(2)370 x(1:2) = __ldca(b(i:j))371 x = __ldcg(b(i:j))372 x = __ldcs(b(i:j))373 x(1:2) = __ldlu(b(i:j))374 x(1:2) = __ldcv(b(i:j))375end376 377! CHECK-LABEL: func.func @_QP__ldxxr2378! CHECK: fir.call @__ldca_r2x2_(%{{.*}}, %{{.*}}) fastmath<contract> : (!fir.ref<!fir.array<2xf16>>, !fir.ref<!fir.array<?xf16>>) -> ()379! CHECK: fir.call @__ldcg_r2x2_(%{{.*}}, %{{.*}}) fastmath<contract> : (!fir.ref<!fir.array<2xf16>>, !fir.ref<!fir.array<?xf16>>) -> ()380! CHECK: fir.call @__ldcs_r2x2_(%{{.*}}, %{{.*}}) fastmath<contract> : (!fir.ref<!fir.array<2xf16>>, !fir.ref<!fir.array<?xf16>>) -> ()381! CHECK: fir.call @__ldlu_r2x2_(%{{.*}}, %{{.*}}) fastmath<contract> : (!fir.ref<!fir.array<2xf16>>, !fir.ref<!fir.array<?xf16>>) -> ()382! CHECK: fir.call @__ldcv_r2x2_(%{{.*}}, %{{.*}}) fastmath<contract> : (!fir.ref<!fir.array<2xf16>>, !fir.ref<!fir.array<?xf16>>) -> ()383 384attributes(global) subroutine __ldXXr8(b)385 real(8), device :: b(*)386 real(8), device :: x(2)387 x(1:2) = __ldca(b(i:j))388 x = __ldcg(b(i:j))389 x = __ldcs(b(i:j))390 x(1:2) = __ldlu(b(i:j))391 x(1:2) = __ldcv(b(i:j))392end393 394! CHECK-LABEL: func.func @_QP__ldxxr8395! CHECK: fir.call @__ldca_r8x2_(%{{.*}}, %{{.*}}) fastmath<contract> : (!fir.ref<!fir.array<2xf64>>, !fir.ref<!fir.array<?xf64>>) -> ()396! CHECK: fir.call @__ldcg_r8x2_(%{{.*}}, %{{.*}}) fastmath<contract> : (!fir.ref<!fir.array<2xf64>>, !fir.ref<!fir.array<?xf64>>) -> ()397! CHECK: fir.call @__ldcs_r8x2_(%{{.*}}, %{{.*}}) fastmath<contract> : (!fir.ref<!fir.array<2xf64>>, !fir.ref<!fir.array<?xf64>>) -> ()398! CHECK: fir.call @__ldlu_r8x2_(%{{.*}}, %{{.*}}) fastmath<contract> : (!fir.ref<!fir.array<2xf64>>, !fir.ref<!fir.array<?xf64>>) -> ()399! CHECK: fir.call @__ldcv_r8x2_(%{{.*}}, %{{.*}}) fastmath<contract> : (!fir.ref<!fir.array<2xf64>>, !fir.ref<!fir.array<?xf64>>) -> ()400 401attributes(device) subroutine testVote()402 integer :: a, ipred, mask403 logical(4) :: pred404 a = all_sync(mask, pred)405 a = any_sync(mask, pred)406 a = ballot_sync(mask, pred)407end subroutine408 409! CHECK-LABEL: func.func @_QPtestvote()410! CHECK: %{{.*}} = nvvm.vote.sync all %{{.*}}, %{{.*}} -> i1411! CHECK: %{{.*}} = nvvm.vote.sync any %{{.*}}, %{{.*}} -> i1412! CHECK: %{{.*}} = nvvm.vote.sync ballot %{{.*}}, %{{.*}} -> i32413 414attributes(global) subroutine test_barrier()415 integer(8), shared :: barrier416 integer(8) :: token417 integer :: count418 call barrier_init(barrier, 256)419 420 token = barrier_arrive(barrier)421 422 token = barrier_arrive(barrier, count)423end subroutine424 425! CHECK-LABEL: func.func @_QPtest_barrier()426 427! CHECK: %[[SHARED:.*]] = cuf.shared_memory i64 {bindc_name = "barrier", uniq_name = "_QFtest_barrierEbarrier"} -> !fir.ref<i64>428! CHECK: %[[DECL_SHARED:.*]]:2 = hlfir.declare %[[SHARED]] {data_attr = #cuf.cuda<shared>, uniq_name = "_QFtest_barrierEbarrier"} : (!fir.ref<i64>) -> (!fir.ref<i64>, !fir.ref<i64>)429! CHECK: %[[COUNT:.*]] = arith.constant 256 : i32430! CHECK: %[[LLVM_PTR:.*]] = fir.convert %[[DECL_SHARED]]#0 : (!fir.ref<i64>) -> !llvm.ptr431! CHECK: %[[SHARED_PTR:.*]] = llvm.addrspacecast %[[LLVM_PTR]] : !llvm.ptr to !llvm.ptr<3>432! CHECK: nvvm.mbarrier.init %[[SHARED_PTR]], %[[COUNT]] : !llvm.ptr<3>, i32433! CHECK: nvvm.fence.proxy {kind = #nvvm.proxy_kind<async.shared>, space = #nvvm.shared_space<cta>}434 435! CHECK: %[[LLVM_PTR:.*]] = fir.convert %[[DECL_SHARED]]#0 : (!fir.ref<i64>) -> !llvm.ptr436! CHECK: %[[SHARED_PTR:.*]] = llvm.addrspacecast %[[LLVM_PTR]] : !llvm.ptr to !llvm.ptr<3>437! CHECK: %{{.*}} = nvvm.mbarrier.arrive %[[SHARED_PTR]] : !llvm.ptr<3> -> i64438 439! CHECK: %[[LLVM_PTR:.*]] = fir.convert %[[DECL_SHARED]]#0 : (!fir.ref<i64>) -> !llvm.ptr440! CHECK: %[[SHARED_PTR:.*]] = llvm.addrspacecast %[[LLVM_PTR]] : !llvm.ptr to !llvm.ptr<3>441! CHECK: %{{.*}} = nvvm.inline_ptx "mbarrier.arrive.expect_tx.release.cta.shared::cta.b64 %{{.*}}, [%{{.*}}], %{{.*}};" ro(%{{.*}}, %{{.*}} : !llvm.ptr<3>, i32) -> i64442 443 444attributes(global) subroutine test_fence()445 call fence_proxy_async()446end subroutine447 448! CHECK-LABEL: func.func @_QPtest_fence()449! CHECK: nvvm.fence.proxy {kind = #nvvm.proxy_kind<async.shared>, space = #nvvm.shared_space<cta>}450 451attributes(global) subroutine test_tma()452 call tma_bulk_commit_group()453 call tma_bulk_wait_group()454end subroutine455 456! CHECK-LABEL: func.func @_QPtest_tma()457! CHECK: nvvm.cp.async.bulk.commit.group458! CHECK: nvvm.cp.async.bulk.wait_group 0459 460attributes(global) subroutine test_bulk_g2s(a)461 real(8), device :: a(*)462 real(8), shared :: tmpa(1024)463 integer(8), shared :: barrier1464 integer(4) :: tx_count465 call tma_bulk_g2s(barrier1, a(j), tmpa, tx_count)466end subroutine467 468! CHECK-LABEL: func.func @_QPtest_bulk_g2s469! CHECK: %[[BARRIER:.*]]:2 = hlfir.declare %4 {data_attr = #cuf.cuda<shared>, uniq_name = "_QFtest_bulk_g2sEbarrier1"} : (!fir.ref<i64>) -> (!fir.ref<i64>, !fir.ref<i64>)470! CHECK: %[[DST:.*]]:2 = hlfir.declare %16(%17) {data_attr = #cuf.cuda<shared>, uniq_name = "_QFtest_bulk_g2sEtmpa"} : (!fir.ref<!fir.array<1024xf64>>, !fir.shape<1>) -> (!fir.ref<!fir.array<1024xf64>>, !fir.ref<!fir.array<1024xf64>>)471! CHECK: %[[COUNT:.*]]:2 = hlfir.declare %19 {data_attr = #cuf.cuda<device>, uniq_name = "_QFtest_bulk_g2sEtx_count"} : (!fir.ref<i32>) -> (!fir.ref<i32>, !fir.ref<i32>) 472! CHECK: %[[SRC:.*]] = hlfir.designate %{{.*}} (%{{.*}}) : (!fir.box<!fir.array<?xf64>>, i64) -> !fir.ref<f64>473! CHECK: %[[COUNT_LOAD:.*]] = fir.load %20#0 : !fir.ref<i32>474! CHECK: %[[BARRIER_PTR:.*]] = fir.convert %[[BARRIER]]#0 : (!fir.ref<i64>) -> !llvm.ptr475! CHECK: %[[BARRIER_3:.*]] = llvm.addrspacecast %[[BARRIER_PTR]] : !llvm.ptr to !llvm.ptr<3>476! CHECK: %[[DST_PTR:.*]] = fir.convert %[[DST]]#0 : (!fir.ref<!fir.array<1024xf64>>) -> !llvm.ptr477! CHECK: %[[DST_7:.*]] = llvm.addrspacecast %[[DST_PTR]] : !llvm.ptr to !llvm.ptr<7>478! CHECK: %[[SRC_PTR:.*]] = fir.convert %[[SRC]] : (!fir.ref<f64>) -> !llvm.ptr479! CHECK: %[[SRC_3:.*]] = llvm.addrspacecast %[[SRC_PTR]] : !llvm.ptr to !llvm.ptr<1>480! CHECK: nvvm.cp.async.bulk.shared.cluster.global %[[DST_7]], %[[SRC_3]], %[[BARRIER_3]], %[[COUNT_LOAD]] : !llvm.ptr<7>, <1>481 482attributes(global) subroutine test_bulk_s2g(a)483 real(8), device :: a(*)484 real(8), shared :: tmpa(1024)485 integer(4) :: tx_count486 call tma_bulk_s2g(tmpa, a(j), tx_count)487end subroutine488 489! CHECK-LABEL: func.func @_QPtest_bulk_s2g490! CHECL: nvvm.cp.async.bulk.global.shared.cta %{{.*}}, %{{.*}}, %{{.*}} : <1>, <3>491! CHECK: nvvm.inline_ptx "cp.async.bulk.commit_group;"492! CHECK: nvvm.cp.async.bulk.wait_group 0493 494attributes(device) subroutine testAtomicCasLoop(aa, n)495 integer :: a496 do while (atomiccas(a, 0, 1) == 1)497 end do498end subroutine499 500! CHECK-LABEL: func.func @_QPtestatomiccasloop501! CHECK: %[[CMP_XCHG:.*]] = llvm.cmpxchg %15, %c0_i32, %c1_i32 acq_rel monotonic : !llvm.ptr, i32502! CHECK: %[[CMP_XCHG_EV:.*]] = llvm.extractvalue %[[CMP_XCHG]][1] : !llvm.struct<(i32, i1)> 503! CHECK: %[[CASTED_CMP_XCHG_EV:.*]] = fir.convert %[[CMP_XCHG_EV]] : (i1) -> i32504! CHECK: %{{.*}} = arith.constant 1 : i32505! CHECK: %19 = arith.cmpi eq, %[[CASTED_CMP_XCHG_EV]], %{{.*}} : i32506 507attributes(global) subroutine test_barrier_try_wait()508 integer :: istat509 integer(8), shared :: barrier1510 integer(8) :: token511 istat = barrier_try_wait(barrier1, token)512end subroutine513 514! CHECK-LABEL: func.func @_QPtest_barrier_try_wait()515! CHECK: scf.while516! CHECK: %{{.*}} = nvvm.inline_ptx "{\0A .reg .pred p;\0A mbarrier.try_wait.shared.b64 p, [%{{.*}}], %{{.*}}, %{{.*}};\0A selp.b32 %{{.*}}, 1, 0, p;\0A}" ro(%{{.*}}, %{{.*}}, %{{.*}} : !llvm.ptr, i64, i32) -> i32517 518attributes(global) subroutine test_barrier_try_wait_sleep()519 integer :: istat520 integer(8), shared :: barrier1521 integer(8) :: token522 integer(4) :: sleep_time523 istat = barrier_try_wait_sleep(barrier1, token, sleep_time)524end subroutine525 526! CHECK-LABEL: func.func @_QPtest_barrier_try_wait_sleep()527! CHECK: %{{.*}} = nvvm.inline_ptx "{\0A .reg .pred p;\0A mbarrier.try_wait.shared.b64 p, [%{{.*}}], %{{.*}}, %{{.*}};\0A selp.b32 %{{.*}}, 1, 0, p;\0A}" ro(%{{.*}}, %{{.*}}, %{{.*}} : !llvm.ptr, i64, i32) -> i32528 529attributes(global) subroutine test_tma_bulk_load_c4(a, n)530 integer(8), shared :: barrier1531 integer, value :: n532 complex(4), device :: r8(n)533 complex(4), shared :: tmp(1024)534 integer(4) :: j, elem_count535 call tma_bulk_load(barrier1, r8(j), tmp, elem_count)536end subroutine537 538! CHECK-LABEL: func.func @_QPtest_tma_bulk_load_c4539! CHECK: %[[BARRIER:.*]]:2 = hlfir.declare %{{.*}} {data_attr = #cuf.cuda<shared>, uniq_name = "_QFtest_tma_bulk_load_c4Ebarrier1"} : (!fir.ref<i64>) -> (!fir.ref<i64>, !fir.ref<i64>)540! CHECK: %[[ELEM_COUNT:.*]]:2 = hlfir.declare %{{.*}} {data_attr = #cuf.cuda<device>, uniq_name = "_QFtest_tma_bulk_load_c4Eelem_count"} : (!fir.ref<i32>) -> (!fir.ref<i32>, !fir.ref<i32>)541! CHECK: %[[COUNT:.*]] = fir.load %[[ELEM_COUNT]]#0 : !fir.ref<i32>542! CHECK: %[[ELEM_SIZE:.*]] = arith.constant 8 : i32543! CHECK: %[[SIZE:.*]] = arith.muli %[[COUNT]], %[[ELEM_SIZE]] : i32544! CHECK: %[[BARRIER_PTR:.*]] = fir.convert %[[BARRIER]]#0 : (!fir.ref<i64>) -> !llvm.ptr545! CHECK: nvvm.inline_ptx "cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%0], [%1], %2, [%3];" ro(%{{.*}}, %{{.*}}, %[[SIZE]], %[[BARRIER_PTR]] : !llvm.ptr, !llvm.ptr, i32, !llvm.ptr)546! CHECK: nvvm.inline_ptx "mbarrier.expect_tx.relaxed.cta.shared::cta.b64 [%0], %1;" ro(%[[BARRIER_PTR]], %[[SIZE]] : !llvm.ptr, i32)547 548attributes(global) subroutine test_tma_bulk_load_c8(a, n)549 integer(8), shared :: barrier1550 integer, value :: n551 complex(8), device :: r8(n)552 complex(8), shared :: tmp(1024)553 integer(4) :: j, elem_count554 call tma_bulk_load(barrier1, r8(j), tmp, elem_count)555end subroutine556 557! CHECK-LABEL: func.func @_QPtest_tma_bulk_load_c8558! CHECK: %[[BARRIER:.*]]:2 = hlfir.declare %{{.*}} {data_attr = #cuf.cuda<shared>, uniq_name = "_QFtest_tma_bulk_load_c8Ebarrier1"} : (!fir.ref<i64>) -> (!fir.ref<i64>, !fir.ref<i64>)559! CHECK: %[[ELEM_COUNT:.*]]:2 = hlfir.declare %{{.*}} {data_attr = #cuf.cuda<device>, uniq_name = "_QFtest_tma_bulk_load_c8Eelem_count"} : (!fir.ref<i32>) -> (!fir.ref<i32>, !fir.ref<i32>)560! CHECK: %[[COUNT:.*]] = fir.load %[[ELEM_COUNT]]#0 : !fir.ref<i32>561! CHECK: %[[ELEM_SIZE:.*]] = arith.constant 16 : i32562! CHECK: %[[SIZE:.*]] = arith.muli %[[COUNT]], %[[ELEM_SIZE]] : i32563! CHECK: %[[BARRIER_PTR:.*]] = fir.convert %[[BARRIER]]#0 : (!fir.ref<i64>) -> !llvm.ptr564! CHECK: nvvm.inline_ptx "cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%0], [%1], %2, [%3];" ro(%{{.*}}, %{{.*}}, %[[SIZE]], %[[BARRIER_PTR]] : !llvm.ptr, !llvm.ptr, i32, !llvm.ptr)565! CHECK: nvvm.inline_ptx "mbarrier.expect_tx.relaxed.cta.shared::cta.b64 [%0], %1;" ro(%[[BARRIER_PTR]], %[[SIZE]] : !llvm.ptr, i32)566 567attributes(global) subroutine test_tma_bulk_load_i4(a, n)568 integer(8), shared :: barrier1569 integer, value :: n570 integer(4), device :: r8(n)571 integer(4), shared :: tmp(1024)572 integer(4) :: j, elem_count573 call tma_bulk_load(barrier1, r8(j), tmp, elem_count)574end subroutine575 576! CHECK-LABEL: func.func @_QPtest_tma_bulk_load_i4577! CHECK: %[[BARRIER:.*]]:2 = hlfir.declare %{{.*}} {data_attr = #cuf.cuda<shared>, uniq_name = "_QFtest_tma_bulk_load_i4Ebarrier1"} : (!fir.ref<i64>) -> (!fir.ref<i64>, !fir.ref<i64>)578! CHECK: %[[ELEM_COUNT:.*]]:2 = hlfir.declare %{{.*}} {data_attr = #cuf.cuda<device>, uniq_name = "_QFtest_tma_bulk_load_i4Eelem_count"} : (!fir.ref<i32>) -> (!fir.ref<i32>, !fir.ref<i32>)579! CHECK: %[[COUNT:.*]] = fir.load %[[ELEM_COUNT]]#0 : !fir.ref<i32>580! CHECK: %[[ELEM_SIZE:.*]] = arith.constant 4 : i32581! CHECK: %[[SIZE:.*]] = arith.muli %[[COUNT]], %[[ELEM_SIZE]] : i32582! CHECK: %[[BARRIER_PTR:.*]] = fir.convert %[[BARRIER]]#0 : (!fir.ref<i64>) -> !llvm.ptr583! CHECK: nvvm.inline_ptx "cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%0], [%1], %2, [%3];" ro(%{{.*}}, %{{.*}}, %[[SIZE]], %[[BARRIER_PTR]] : !llvm.ptr, !llvm.ptr, i32, !llvm.ptr)584! CHECK: nvvm.inline_ptx "mbarrier.expect_tx.relaxed.cta.shared::cta.b64 [%0], %1;" ro(%[[BARRIER_PTR]], %[[SIZE]] : !llvm.ptr, i32)585 586attributes(global) subroutine test_tma_bulk_load_i8(a, n)587 integer(8), shared :: barrier1588 integer, value :: n589 integer(8), device :: r8(n)590 integer(8), shared :: tmp(1024)591 integer(4) :: j, elem_count592 call tma_bulk_load(barrier1, r8(j), tmp, elem_count)593end subroutine594 595! CHECK-LABEL: func.func @_QPtest_tma_bulk_load_i8596! CHECK: %[[BARRIER:.*]]:2 = hlfir.declare %{{.*}} {data_attr = #cuf.cuda<shared>, uniq_name = "_QFtest_tma_bulk_load_i8Ebarrier1"} : (!fir.ref<i64>) -> (!fir.ref<i64>, !fir.ref<i64>)597! CHECK: %[[ELEM_COUNT:.*]]:2 = hlfir.declare %{{.*}} {data_attr = #cuf.cuda<device>, uniq_name = "_QFtest_tma_bulk_load_i8Eelem_count"} : (!fir.ref<i32>) -> (!fir.ref<i32>, !fir.ref<i32>)598! CHECK: %[[COUNT:.*]] = fir.load %[[ELEM_COUNT]]#0 : !fir.ref<i32>599! CHECK: %[[ELEM_SIZE:.*]] = arith.constant 8 : i32600! CHECK: %[[SIZE:.*]] = arith.muli %[[COUNT]], %[[ELEM_SIZE]] : i32601! CHECK: %[[BARRIER_PTR:.*]] = fir.convert %[[BARRIER]]#0 : (!fir.ref<i64>) -> !llvm.ptr602! CHECK: nvvm.inline_ptx "cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%0], [%1], %2, [%3];" ro(%{{.*}}, %{{.*}}, %[[SIZE]], %[[BARRIER_PTR]] : !llvm.ptr, !llvm.ptr, i32, !llvm.ptr)603! CHECK: nvvm.inline_ptx "mbarrier.expect_tx.relaxed.cta.shared::cta.b64 [%0], %1;" ro(%[[BARRIER_PTR]], %[[SIZE]] : !llvm.ptr, i32)604 605attributes(global) subroutine test_tma_bulk_load_r2(a, n)606 integer(8), shared :: barrier1607 integer, value :: n608 real(2), device :: r8(n)609 real(2), shared :: tmp(1024)610 integer(4) :: j, elem_count611 call tma_bulk_load(barrier1, r8(j), tmp, elem_count)612end subroutine613 614! CHECK-LABEL: func.func @_QPtest_tma_bulk_load_r2615! CHECK: %[[BARRIER:.*]]:2 = hlfir.declare %{{.*}} {data_attr = #cuf.cuda<shared>, uniq_name = "_QFtest_tma_bulk_load_r2Ebarrier1"} : (!fir.ref<i64>) -> (!fir.ref<i64>, !fir.ref<i64>)616! CHECK: %[[ELEM_COUNT:.*]]:2 = hlfir.declare %{{.*}} {data_attr = #cuf.cuda<device>, uniq_name = "_QFtest_tma_bulk_load_r2Eelem_count"} : (!fir.ref<i32>) -> (!fir.ref<i32>, !fir.ref<i32>)617! CHECK: %[[COUNT:.*]] = fir.load %[[ELEM_COUNT]]#0 : !fir.ref<i32>618! CHECK: %[[ELEM_SIZE:.*]] = arith.constant 2 : i32619! CHECK: %[[SIZE:.*]] = arith.muli %[[COUNT]], %[[ELEM_SIZE]] : i32620! CHECK: %[[BARRIER_PTR:.*]] = fir.convert %[[BARRIER]]#0 : (!fir.ref<i64>) -> !llvm.ptr621! CHECK: nvvm.inline_ptx "cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%0], [%1], %2, [%3];" ro(%{{.*}}, %{{.*}}, %[[SIZE]], %[[BARRIER_PTR]] : !llvm.ptr, !llvm.ptr, i32, !llvm.ptr)622! CHECK: nvvm.inline_ptx "mbarrier.expect_tx.relaxed.cta.shared::cta.b64 [%0], %1;" ro(%[[BARRIER_PTR]], %[[SIZE]] : !llvm.ptr, i32)623 624attributes(global) subroutine test_tma_bulk_load_r4(a, n)625 integer(8), shared :: barrier1626 integer, value :: n627 real(4), device :: r8(n)628 real(4), shared :: tmp(1024)629 integer(4) :: j, elem_count630 call tma_bulk_load(barrier1, r8(j), tmp, elem_count)631end subroutine632 633! CHECK-LABEL: func.func @_QPtest_tma_bulk_load_r4634! CHECK: %[[BARRIER:.*]]:2 = hlfir.declare %{{.*}} {data_attr = #cuf.cuda<shared>, uniq_name = "_QFtest_tma_bulk_load_r4Ebarrier1"} : (!fir.ref<i64>) -> (!fir.ref<i64>, !fir.ref<i64>)635! CHECK: %[[ELEM_COUNT:.*]]:2 = hlfir.declare %{{.*}} {data_attr = #cuf.cuda<device>, uniq_name = "_QFtest_tma_bulk_load_r4Eelem_count"} : (!fir.ref<i32>) -> (!fir.ref<i32>, !fir.ref<i32>)636! CHECK: %[[COUNT:.*]] = fir.load %[[ELEM_COUNT]]#0 : !fir.ref<i32>637! CHECK: %[[ELEM_SIZE:.*]] = arith.constant 4 : i32638! CHECK: %[[SIZE:.*]] = arith.muli %[[COUNT]], %[[ELEM_SIZE]] : i32639! CHECK: %[[BARRIER_PTR:.*]] = fir.convert %[[BARRIER]]#0 : (!fir.ref<i64>) -> !llvm.ptr640! CHECK: nvvm.inline_ptx "cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%0], [%1], %2, [%3];" ro(%{{.*}}, %{{.*}}, %[[SIZE]], %[[BARRIER_PTR]] : !llvm.ptr, !llvm.ptr, i32, !llvm.ptr)641! CHECK: nvvm.inline_ptx "mbarrier.expect_tx.relaxed.cta.shared::cta.b64 [%0], %1;" ro(%[[BARRIER_PTR]], %[[SIZE]] : !llvm.ptr, i32)642 643attributes(global) subroutine test_tma_bulk_load_r8(a, n)644 integer(8), shared :: barrier1645 integer, value :: n646 real(8), device :: r8(n)647 real(8), shared :: tmp(1024)648 integer(4) :: j, elem_count649 call tma_bulk_load(barrier1, r8(j), tmp, elem_count)650end subroutine651 652! CHECK-LABEL: func.func @_QPtest_tma_bulk_load_r8653! CHECK: %[[BARRIER:.*]]:2 = hlfir.declare %{{.*}} {data_attr = #cuf.cuda<shared>, uniq_name = "_QFtest_tma_bulk_load_r8Ebarrier1"} : (!fir.ref<i64>) -> (!fir.ref<i64>, !fir.ref<i64>)654! CHECK: %[[ELEM_COUNT:.*]]:2 = hlfir.declare %{{.*}} {data_attr = #cuf.cuda<device>, uniq_name = "_QFtest_tma_bulk_load_r8Eelem_count"} : (!fir.ref<i32>) -> (!fir.ref<i32>, !fir.ref<i32>)655! CHECK: %[[COUNT:.*]] = fir.load %[[ELEM_COUNT]]#0 : !fir.ref<i32>656! CHECK: %[[ELEM_SIZE:.*]] = arith.constant 8 : i32657! CHECK: %[[SIZE:.*]] = arith.muli %[[COUNT]], %[[ELEM_SIZE]] : i32658! CHECK: %[[BARRIER_PTR:.*]] = fir.convert %[[BARRIER]]#0 : (!fir.ref<i64>) -> !llvm.ptr659! CHECK: nvvm.inline_ptx "cp.async.bulk.shared::cluster.global.mbarrier::complete_tx::bytes [%0], [%1], %2, [%3];" ro(%{{.*}}, %{{.*}}, %[[SIZE]], %[[BARRIER_PTR]] : !llvm.ptr, !llvm.ptr, i32, !llvm.ptr)660! CHECK: nvvm.inline_ptx "mbarrier.expect_tx.relaxed.cta.shared::cta.b64 [%0], %1;" ro(%[[BARRIER_PTR]], %[[SIZE]] : !llvm.ptr, i32)661 662attributes(global) subroutine test_tma_bulk_store_c4(c, n)663 integer, value :: n664 complex(4), device :: c(n)665 complex(4), shared :: tmpa(1024)666 integer(4) :: j, elem_count667 call tma_bulk_store(tmpa, c(j), elem_count)668end subroutine669 670! CHECK-LABEL: func.func @_QPtest_tma_bulk_store_c4671! CHECK: nvvm.cp.async.bulk.global.shared.cta %{{.*}}, %{{.*}}, %{{.*}} : <1>, <3>672! CHECK: nvvm.inline_ptx "cp.async.bulk.commit_group;"673! CHECK: nvvm.cp.async.bulk.wait_group 0674 675attributes(global) subroutine test_tma_bulk_store_c8(c, n)676 integer, value :: n677 complex(8), device :: c(n)678 complex(8), shared :: tmpa(1024)679 integer(4) :: j, elem_count680 call tma_bulk_store(tmpa, c(j), elem_count)681end subroutine682 683! CHECK-LABEL: func.func @_QPtest_tma_bulk_store_c8684! CHECK: nvvm.cp.async.bulk.global.shared.cta %{{.*}}, %{{.*}}, %{{.*}} : <1>, <3>685! CHECK: nvvm.inline_ptx "cp.async.bulk.commit_group;"686! CHECK: nvvm.cp.async.bulk.wait_group 0687 688attributes(global) subroutine test_tma_bulk_store_i4(c, n)689 integer, value :: n690 integer(4), device :: c(n)691 integer(4), shared :: tmpa(1024)692 integer(4) :: j, elem_count693 call tma_bulk_store(tmpa, c(j), elem_count)694end subroutine695 696! CHECK-LABEL: func.func @_QPtest_tma_bulk_store_i4697! CHECK: nvvm.cp.async.bulk.global.shared.cta %{{.*}}, %{{.*}}, %{{.*}} : <1>, <3>698! CHECK: nvvm.inline_ptx "cp.async.bulk.commit_group;"699! CHECK: nvvm.cp.async.bulk.wait_group 0700 701attributes(global) subroutine test_tma_bulk_store_i8(c, n)702 integer, value :: n703 integer(8), device :: c(n)704 integer(8), shared :: tmpa(1024)705 integer(4) :: j, elem_count706 call tma_bulk_store(tmpa, c(j), elem_count)707end subroutine708 709! CHECK-LABEL: func.func @_QPtest_tma_bulk_store_i8710! CHECK: nvvm.cp.async.bulk.global.shared.cta %{{.*}}, %{{.*}}, %{{.*}} : <1>, <3>711! CHECK: nvvm.inline_ptx "cp.async.bulk.commit_group;"712! CHECK: nvvm.cp.async.bulk.wait_group 0713 714 715attributes(global) subroutine test_tma_bulk_store_r2(c, n)716 integer, value :: n717 real(2), device :: c(n)718 real(2), shared :: tmpa(1024)719 integer(4) :: j, elem_count720 call tma_bulk_store(tmpa, c(j), elem_count)721end subroutine722 723! CHECK-LABEL: func.func @_QPtest_tma_bulk_store_r2724! CHECK: nvvm.cp.async.bulk.global.shared.cta %{{.*}}, %{{.*}}, %{{.*}} : <1>, <3>725! CHECK: nvvm.inline_ptx "cp.async.bulk.commit_group;"726! CHECK: nvvm.cp.async.bulk.wait_group 0727 728attributes(global) subroutine test_tma_bulk_store_r4(c, n)729 integer, value :: n730 real(4), device :: c(n)731 real(4), shared :: tmpa(1024)732 integer(4) :: j, elem_count733 call tma_bulk_store(tmpa, c(j), elem_count)734end subroutine735 736! CHECK-LABEL: func.func @_QPtest_tma_bulk_store_r4737! CHECK: nvvm.cp.async.bulk.global.shared.cta %{{.*}}, %{{.*}}, %{{.*}} : <1>, <3>738! CHECK: nvvm.inline_ptx "cp.async.bulk.commit_group;"739! CHECK: nvvm.cp.async.bulk.wait_group 0740 741attributes(global) subroutine test_tma_bulk_store_r8(c, n)742 integer, value :: n743 real(8), device :: c(n)744 real(8), shared :: tmpa(1024)745 integer(4) :: j, elem_count746 call tma_bulk_store(tmpa, c(j), elem_count)747end subroutine748 749! CHECK-LABEL: func.func @_QPtest_tma_bulk_store_r8750! CHECK: nvvm.cp.async.bulk.global.shared.cta %{{.*}}, %{{.*}}, %{{.*}} : <1>, <3>751! CHECK: nvvm.inline_ptx "cp.async.bulk.commit_group;"752! CHECK: nvvm.cp.async.bulk.wait_group 0753