42 lines · plain
1// RUN: mlir-translate -mlir-to-llvmir %s | FileCheck %s --check-prefix=CHECK-LLVM2 3// CHECK-LABEL: @llvm_nvvm_tcgen05_alloc4llvm.func @llvm_nvvm_tcgen05_alloc(%addr : !llvm.ptr, %ncols : i32) {5 // CHECK-LLVM: call void @llvm.nvvm.tcgen05.alloc.cg1(ptr %{{.*}}, i32 %{{.*}})6 nvvm.tcgen05.alloc %addr, %ncols : !llvm.ptr, i327 8 // CHECK-LLVM: call void @llvm.nvvm.tcgen05.alloc.cg2(ptr %{{.*}}, i32 %{{.*}})9 nvvm.tcgen05.alloc %addr, %ncols {group = #nvvm.cta_group<cta_2>} : !llvm.ptr, i3210 llvm.return11}12 13// CHECK-LABEL: @llvm_nvvm_tcgen05_alloc_shared14llvm.func @llvm_nvvm_tcgen05_alloc_shared(%addr : !llvm.ptr<3>, %ncols : i32) {15 // CHECK-LLVM: call void @llvm.nvvm.tcgen05.alloc.shared.cg1(ptr addrspace(3) %{{.*}}, i32 %{{.*}})16 nvvm.tcgen05.alloc %addr, %ncols : !llvm.ptr<3>, i3217 18 // CHECK-LLVM: call void @llvm.nvvm.tcgen05.alloc.shared.cg2(ptr addrspace(3) %{{.*}}, i32 %{{.*}})19 nvvm.tcgen05.alloc %addr, %ncols {group = #nvvm.cta_group<cta_2>} : !llvm.ptr<3>, i3220 llvm.return21}22 23// CHECK-LABEL: @llvm_nvvm_tcgen05_dealloc24llvm.func @llvm_nvvm_tcgen05_dealloc(%addr : !llvm.ptr<6>, %ncols : i32) {25 // CHECK-LLVM: call void @llvm.nvvm.tcgen05.dealloc.cg1(ptr addrspace(6) %{{.*}}, i32 %{{.*}})26 nvvm.tcgen05.dealloc %addr, %ncols : !llvm.ptr<6>, i3227 28 // CHECK-LLVM: call void @llvm.nvvm.tcgen05.dealloc.cg2(ptr addrspace(6) %{{.*}}, i32 %{{.*}})29 nvvm.tcgen05.dealloc %addr, %ncols {group = #nvvm.cta_group<cta_2>} : !llvm.ptr<6>, i3230 llvm.return31}32 33// CHECK-LABEL: @llvm_nvvm_tcgen05_relinquish_alloc_permit34llvm.func @llvm_nvvm_tcgen05_relinquish_alloc_permit() {35 // CHECK-LLVM: call void @llvm.nvvm.tcgen05.relinq.alloc.permit.cg1()36 nvvm.tcgen05.relinquish_alloc_permit37 38 // CHECK-LLVM: call void @llvm.nvvm.tcgen05.relinq.alloc.permit.cg2()39 nvvm.tcgen05.relinquish_alloc_permit {group = #nvvm.cta_group<cta_2>}40 llvm.return41}42