115 lines · plain
1// RUN: mlir-translate -verify-diagnostics -split-input-file -mlir-to-llvmir %s2 3// -----4 5llvm.func @mbarrier_arrive_ret_check(%barrier: !llvm.ptr<7>) {6 // expected-error @below {{mbarrier in shared_cluster space cannot return any value}}7 %0 = nvvm.mbarrier.arrive %barrier : !llvm.ptr<7> -> i648 llvm.return9}10 11// -----12 13llvm.func @mbarrier_arrive_invalid_scope(%barrier: !llvm.ptr<7>) {14 // expected-error @below {{mbarrier scope must be either CTA or Cluster}}15 %0 = nvvm.mbarrier.arrive %barrier {scope = #nvvm.mem_scope<gpu>} : !llvm.ptr<7> -> i6416 llvm.return17}18 19// -----20 21llvm.func @mbarrier_arrive_drop_ret_check(%barrier: !llvm.ptr<7>) {22 // expected-error @below {{mbarrier in shared_cluster space cannot return any value}}23 %0 = nvvm.mbarrier.arrive_drop %barrier : !llvm.ptr<7> -> i6424 llvm.return25}26 27// -----28 29llvm.func @mbarrier_arrive_drop_invalid_scope(%barrier: !llvm.ptr<7>) {30 // expected-error @below {{mbarrier scope must be either CTA or Cluster}}31 %0 = nvvm.mbarrier.arrive_drop %barrier {scope = #nvvm.mem_scope<gpu>} : !llvm.ptr<7> -> i6432 llvm.return33}34 35// -----36 37llvm.func @mbarrier_expect_tx_scope(%barrier: !llvm.ptr<7>, %tx_count: i32) {38 // expected-error @below {{mbarrier scope must be either CTA or Cluster}}39 nvvm.mbarrier.expect_tx %barrier, %tx_count {scope = #nvvm.mem_scope<gpu>} : !llvm.ptr<7>, i3240 llvm.return41}42 43// -----44 45llvm.func @mbarrier_complete_tx_scope(%barrier: !llvm.ptr<3>, %tx_count: i32) {46 // expected-error @below {{mbarrier scope must be either CTA or Cluster}}47 nvvm.mbarrier.complete_tx %barrier, %tx_count {scope = #nvvm.mem_scope<sys>} : !llvm.ptr<3>, i3248 llvm.return49}50 51// -----52 53llvm.func @mbarrier_arr_expect_tx(%barrier: !llvm.ptr<3>, %tx_count: i32) {54 // expected-error @below {{mbarrier scope must be either CTA or Cluster}}55 %1 = nvvm.mbarrier.arrive.expect_tx %barrier, %tx_count {scope = #nvvm.mem_scope<gpu>} : !llvm.ptr<3>, i32 -> i6456 llvm.return57}58 59// -----60 61llvm.func @mbarrier_arr_expect_tx_cluster(%barrier: !llvm.ptr<7>, %tx_count: i32) {62 // expected-error @below {{mbarrier in shared_cluster space cannot return any value}}63 %1 = nvvm.mbarrier.arrive.expect_tx %barrier, %tx_count {scope = #nvvm.mem_scope<cta>} : !llvm.ptr<7>, i32 -> i6464 llvm.return65}66 67// -----68 69llvm.func @init_mbarrier_arrive_expect_tx_asm_ret(%barrier : !llvm.ptr<3>, %txcount : i32, %pred : i1) {70 // expected-error @below {{return-value is not supported when using predicate}}71 %1 = nvvm.mbarrier.arrive.expect_tx %barrier, %txcount, predicate = %pred : !llvm.ptr<3>, i32, i1 -> i64 72 llvm.return73}74 75// -----76 77llvm.func @init_mbarrier_arrive_expect_tx_asm_relaxed(%barrier : !llvm.ptr<3>, %txcount : i32, %pred : i1) {78 // expected-error @below {{mbarrier with relaxed semantics is not supported when using predicate}}79 nvvm.mbarrier.arrive.expect_tx %barrier, %txcount, predicate = %pred {relaxed = true} : !llvm.ptr<3>, i32, i180 llvm.return81}82 83// -----84 85llvm.func @init_mbarrier_arrive_expect_tx_asm_cta(%barrier : !llvm.ptr<3>, %txcount : i32, %pred : i1) {86 // expected-error @below {{mbarrier scope must be CTA when using predicate}}87 nvvm.mbarrier.arrive.expect_tx %barrier, %txcount, predicate = %pred {scope = #nvvm.mem_scope<cluster>} : !llvm.ptr<3>, i32, i188 llvm.return89}90 91// -----92 93llvm.func @init_mbarrier_arrive_expect_tx_asm_cluster(%barrier : !llvm.ptr<7>, %txcount : i32, %pred : i1) {94 // expected-error @below {{mbarrier in shared_cluster space is not supported when using predicate}}95 nvvm.mbarrier.arrive.expect_tx %barrier, %txcount, predicate = %pred : !llvm.ptr<7>, i32, i196 llvm.return97}98 99// -----100 101llvm.func @mbarrier_arr_drop_expect_tx(%barrier: !llvm.ptr<3>, %tx_count: i32) {102 // expected-error @below {{mbarrier scope must be either CTA or Cluster}}103 %1 = nvvm.mbarrier.arrive_drop.expect_tx %barrier, %tx_count {scope = #nvvm.mem_scope<gpu>} : !llvm.ptr<3>, i32 -> i64104 llvm.return105}106 107// -----108 109llvm.func @mbarrier_arr_drop_expect_tx_cluster(%barrier: !llvm.ptr<7>, %tx_count: i32) {110 // expected-error @below {{mbarrier in shared_cluster space cannot return any value}}111 %1 = nvvm.mbarrier.arrive_drop.expect_tx %barrier, %tx_count {scope = #nvvm.mem_scope<cta>} : !llvm.ptr<7>, i32 -> i64112 llvm.return113}114 115