1// RUN: mlir-translate -verify-diagnostics -split-input-file -mlir-to-llvmir %s 2 3// ----- 4 5llvm.func @kernel_func(%numberOfThreads : i32) { 6 // expected-error @below {{'nvvm.barrier' op barrier id is missing, it should be set between 0 to 15}} 7 nvvm.barrier number_of_threads = %numberOfThreads 8} 9 10// ----- 11 12// expected-error @below {{'"nvvm.minctasm"' attribute must be integer constant}} 13llvm.func @kernel_func() attributes {nvvm.kernel, nvvm.minctasm = "foo"} { 14 llvm.return 15} 16 17// ----- 18 19// expected-error @below {{'"nvvm.maxnreg"' attribute must be integer constant}} 20llvm.func @kernel_func() attributes {nvvm.kernel, nvvm.maxnreg = "boo"} { 21 llvm.return 22} 23 24// ----- 25 26// expected-error @below {{'"nvvm.reqntid"' attribute must be integer array with maximum 3 index}} 27llvm.func @kernel_func() attributes {nvvm.kernel, nvvm.reqntid = array<i32: 3, 4, 5, 6>} { 28 llvm.return 29} 30 31// ----- 32 33// expected-error @below {{'"nvvm.maxntid"' attribute must be integer array with maximum 3 index}} 34llvm.func @kernel_func() attributes {nvvm.kernel, nvvm.maxntid = array<i32: 3, 4, 5, 6>} { 35 llvm.return 36} 37 38// ----- 39 40llvm.func @nvvm_fence_proxy_acquire(%addr : !llvm.ptr, %size : i32) { 41 // expected-error @below {{'nvvm.fence.proxy.acquire' op uni-directional proxies only support generic for from_proxy attribute}} 42 nvvm.fence.proxy.acquire #nvvm.mem_scope<cta> %addr, %size from_proxy=#nvvm.proxy_kind<tensormap> to_proxy=#nvvm.proxy_kind<generic> 43 llvm.return 44} 45 46// ----- 47 48llvm.func @nvvm_fence_proxy_release() { 49 // expected-error @below {{'nvvm.fence.proxy.release' op uni-directional proxies only support generic for from_proxy attribute}} 50 nvvm.fence.proxy.release #nvvm.mem_scope<cta> from_proxy=#nvvm.proxy_kind<tensormap> to_proxy=#nvvm.proxy_kind<generic> 51 llvm.return 52} 53 54// ----- 55 56llvm.func @nvvm_fence_proxy_acquire(%addr : !llvm.ptr, %size : i32) { 57 // expected-error @below {{'nvvm.fence.proxy.acquire' op uni-directional proxies only support tensormap for to_proxy attribute}} 58 nvvm.fence.proxy.acquire #nvvm.mem_scope<cta> %addr, %size from_proxy=#nvvm.proxy_kind<generic> to_proxy=#nvvm.proxy_kind<generic> 59 llvm.return 60} 61 62// ----- 63 64llvm.func @nvvm_fence_proxy_release() { 65 // expected-error @below {{'nvvm.fence.proxy.release' op uni-directional proxies only support tensormap for to_proxy attribute}} 66 nvvm.fence.proxy.release #nvvm.mem_scope<cta> from_proxy=#nvvm.proxy_kind<generic> to_proxy=#nvvm.proxy_kind<generic> 67 llvm.return 68} 69 70// ----- 71 72llvm.func @tma_prefetch_0d(%tma_desc : !llvm.ptr, %d0 : i32, %ch : i64) { 73 // expected-error @below {{expects coordinates between 1 to 5 dimension}} 74 nvvm.cp.async.bulk.tensor.prefetch %tma_desc, box[] : !llvm.ptr 75 llvm.return 76} 77 78// ----- 79 80llvm.func @tma_prefetch_2d_im2col(%tma_desc : !llvm.ptr, %d0 : i32, %d1 : i32, %off0 : i16, %ch : i64) { 81 // expected-error @below {{to use im2col mode, the tensor has to be at least 3-dimensional}} 82 nvvm.cp.async.bulk.tensor.prefetch %tma_desc, box[%d0, %d1] im2col[%off0] l2_cache_hint = %ch : !llvm.ptr 83 llvm.return 84} 85 86// ----- 87 88llvm.func @tma_prefetch_5d_im2col(%tma_desc : !llvm.ptr, %d0 : i32, %d1 : i32, %d2 : i32, %d3 : i32, %d4 : i32, %off0 : i16, %off1 : i16, %off2 : i16, %ch : i64) { 89 // expected-error @below {{im2col offsets must be 2 less than number of coordinates}} 90 nvvm.cp.async.bulk.tensor.prefetch %tma_desc, box[%d0, %d1, %d2, %d3, %d4] im2col[%off0, %off1] : !llvm.ptr 91 llvm.return 92} 93 94// ----- 95 96llvm.func @tma_reduce_0d(%src : !llvm.ptr<3>, %tma_desc : !llvm.ptr, %ch : i64) { 97 // expected-error @below {{expects coordinates between 1 to 5 dimension}} 98 nvvm.cp.async.bulk.tensor.reduce %tma_desc, %src, box[] {redKind = #nvvm.tma_redux_kind<add>}: !llvm.ptr, !llvm.ptr<3> 99 llvm.return 100} 101 102// ----- 103 104llvm.func @tma_reduce_2d_im2col(%src : !llvm.ptr<3>, %tma_desc : !llvm.ptr, %d0 : i32, %d1 : i32, %ch : i64) { 105 // expected-error @below {{to use im2col mode, the tensor has to be at least 3-dimensional}} 106 nvvm.cp.async.bulk.tensor.reduce %tma_desc, %src, box[%d0, %d1] {redKind = #nvvm.tma_redux_kind<and>, mode = #nvvm.tma_store_mode<im2col>}: !llvm.ptr, !llvm.ptr<3> 107 llvm.return 108} 109 110// ----- 111 112llvm.func @convert_float_to_tf32_rna_relu(%src : f32) -> i32 { 113 // expected-error @below {{Relu not supported with rna rounding mode.}} 114 %res = nvvm.cvt.float.to.tf32 %src {rnd = #nvvm.fp_rnd_mode<rna>, relu=true} 115 llvm.return %res : i32 116} 117 118// ----- 119 120llvm.func @convert_float_to_tf32_rn_sf(%src : f32) -> i32 { 121 // expected-error @below {{Saturation mode not supported with rn/rz rounding modes.}} 122 %res = nvvm.cvt.float.to.tf32 %src {rnd = #nvvm.fp_rnd_mode<rn>, sat = #nvvm.sat_mode<satfinite>} 123 llvm.return %res : i32 124} 125 126// ----- 127 128llvm.func @convert_float_to_tf32_rz_sf(%src : f32) -> i32 { 129 // expected-error @below {{Saturation mode not supported with rn/rz rounding modes.}} 130 %res = nvvm.cvt.float.to.tf32 %src {rnd = #nvvm.fp_rnd_mode<rz>, sat = #nvvm.sat_mode<satfinite>} 131 llvm.return %res : i32 132} 133 134// ----- 135 136llvm.func @convert_float_to_tf32_no_rnd_mode(%src : f32) -> i32 { 137 // expected-error @below {{Only {rn,rz,rna} rounding modes supported for CvtFloatToTF32Op.}} 138 %res = nvvm.cvt.float.to.tf32 %src 139 llvm.return %res : i32 140} 141