627 lines · plain
1// RUN: mlir-translate -verify-diagnostics -split-input-file -mlir-to-llvmir %s2 3llvm.func @pmevent_no_id() {4 // expected-error @below {{either `id` or `mask` must be set}}5 nvvm.pmevent 6}7 8// -----9 10llvm.func @pmevent_bigger15() {11 // expected-error @below {{`id` must be between 0 and 15}}12 nvvm.pmevent id = 14113}14 15// -----16 17llvm.func @pmevent_many_ids() {18 // expected-error @below {{`id` and `mask` cannot be set at the same time}}19 nvvm.pmevent id = 1 mask = 120}21 22// -----23 24llvm.func @kernel_func(%numberOfThreads : i32) {25 // expected-error @below {{'nvvm.barrier' op barrier id is missing, it should be set between 0 to 15}}26 nvvm.barrier number_of_threads = %numberOfThreads27}28 29// -----30 31// expected-error @below {{'"nvvm.minctasm"' attribute must be integer constant}}32llvm.func @kernel_func() attributes {nvvm.kernel, nvvm.minctasm = "foo"} {33 llvm.return34}35 36// -----37 38// expected-error @below {{'"nvvm.maxnreg"' attribute must be integer constant}}39llvm.func @kernel_func() attributes {nvvm.kernel, nvvm.maxnreg = "boo"} {40 llvm.return41}42 43// -----44 45// expected-error @below {{'"nvvm.reqntid"' attribute must be integer array with maximum 3 index}}46llvm.func @kernel_func() attributes {nvvm.kernel, nvvm.reqntid = array<i32: 3, 4, 5, 6>} {47 llvm.return48}49 50// -----51 52// expected-error @below {{'"nvvm.maxntid"' attribute must be integer array with maximum 3 index}}53llvm.func @kernel_func() attributes {nvvm.kernel, nvvm.maxntid = array<i32: 3, 4, 5, 6>} {54 llvm.return55}56 57// -----58 59// expected-error @below {{'"nvvm.blocksareclusters"' attribute must be used along with 'nvvm.reqntid' and 'nvvm.cluster_dim'}}60llvm.func @kernel_func() attributes {nvvm.kernel, nvvm.blocksareclusters,61 nvvm.cluster_dim = array<i32: 3, 5, 7>} {62 llvm.return63}64 65// -----66 67// expected-error @below {{'"nvvm.blocksareclusters"' attribute must be used along with 'nvvm.reqntid' and 'nvvm.cluster_dim'}}68llvm.func @kernel_func() attributes {nvvm.kernel, nvvm.blocksareclusters,69 nvvm.reqntid = array<i32: 1, 23, 32>} {70 llvm.return71}72 73// -----74 75llvm.func @nvvm_fence_proxy_acquire(%addr : !llvm.ptr, %size : i32) {76 // expected-error @below {{'nvvm.fence.proxy.acquire' op uni-directional proxies only support generic for from_proxy attribute}}77 nvvm.fence.proxy.acquire #nvvm.mem_scope<cta> %addr, %size from_proxy=#nvvm.proxy_kind<tensormap> to_proxy=#nvvm.proxy_kind<generic>78 llvm.return79}80 81// -----82 83llvm.func @nvvm_fence_proxy_release() {84 // expected-error @below {{'nvvm.fence.proxy.release' op uni-directional proxies only support generic for from_proxy attribute}}85 nvvm.fence.proxy.release #nvvm.mem_scope<cta> from_proxy=#nvvm.proxy_kind<tensormap> to_proxy=#nvvm.proxy_kind<generic>86 llvm.return87}88 89// -----90 91llvm.func @nvvm_fence_proxy_acquire(%addr : !llvm.ptr, %size : i32) {92 // expected-error @below {{'nvvm.fence.proxy.acquire' op uni-directional proxies only support tensormap for to_proxy attribute}}93 nvvm.fence.proxy.acquire #nvvm.mem_scope<cta> %addr, %size from_proxy=#nvvm.proxy_kind<generic> to_proxy=#nvvm.proxy_kind<generic>94 llvm.return95}96 97// -----98 99llvm.func @nvvm_fence_proxy_release() {100 // expected-error @below {{'nvvm.fence.proxy.release' op uni-directional proxies only support tensormap for to_proxy attribute}}101 nvvm.fence.proxy.release #nvvm.mem_scope<cta> from_proxy=#nvvm.proxy_kind<generic> to_proxy=#nvvm.proxy_kind<generic>102 llvm.return103}104 105// -----106 107llvm.func @convert_float_to_tf32_rna_relu(%src : f32) -> i32 {108 // expected-error @below {{Relu not supported with rna rounding mode.}}109 %res = nvvm.convert.float.to.tf32 %src {rnd = #nvvm.fp_rnd_mode<rna>, relu=true}110 llvm.return %res : i32111}112 113// -----114 115llvm.func @convert_float_to_tf32_no_rnd_mode(%src : f32) -> i32 {116 // expected-error @below {{Only {rn,rz,rna} rounding modes supported for ConvertFloatToTF32Op.}}117 %res = nvvm.convert.float.to.tf32 %src118 llvm.return %res : i32119}120 121// -----122 123llvm.func @nvvm_st_bulk_initval_nonzero(%addr : !llvm.ptr, %size : i64) {124 // expected-error @below {{only 0 is supported for initVal, got 1}}125 nvvm.st.bulk %addr, size = %size, init = 1 : !llvm.ptr126 llvm.return127}128 129// -----130 131llvm.func @nvvm_tcgen05_cp_128x256b_mc(%taddr : !llvm.ptr<6>, %smem_desc : i64) {132 // expected-error @below {{Invalid multicast type for tcgen05.cp Op}}133 nvvm.tcgen05.cp %taddr, %smem_desc {shape = #nvvm.tcgen05_cp_shape<shape_128x256b>, multicast = #nvvm.tcgen05_cp_multicast<warpx2_02_13>}134 llvm.return135}136 137// -----138 139llvm.func @nvvm_tcgen05_cp_32x128b_wx2(%taddr : !llvm.ptr<6>, %smem_desc : i64) {140 // expected-error @below {{Shape 32x128b requires multicast warpx4 for tcgen05.cp Op}}141 nvvm.tcgen05.cp %taddr, %smem_desc {142 shape = #nvvm.tcgen05_cp_shape<shape_32x128b>,143 multicast = #nvvm.tcgen05_cp_multicast<warpx2_01_23>144 }145 llvm.return146}147 148// -----149 150llvm.func @nvvm_tcgen05_cp_64x128b(%taddr : !llvm.ptr<6>, %smem_desc : i64) {151 // expected-error @below {{Shape 64x128b requires multicast warpx2_01_23 or warpx2_02_13 for tcgen05.cp Op}}152 nvvm.tcgen05.cp %taddr, %smem_desc {153 shape = #nvvm.tcgen05_cp_shape<shape_64x128b>,154 multicast = #nvvm.tcgen05_cp_multicast<warpx4>155 }156 llvm.return157}158 159// -----160 161llvm.func @nvvm_match_sync_all(%val32: i32, %thread_mask: i32) {162 // expected-error @below {{match.sync 'all' returns a two element struct with first element as i32 and second element as i1}}163 %0 = nvvm.match.sync all %thread_mask, %val32 : i32 -> !llvm.struct<(i32, i8)>164 llvm.return165}166 167// -----168 169llvm.func @nvvm_match_sync_any(%val32: i32, %thread_mask: i32) {170 // expected-error @below {{match.sync 'any' returns an i32}}171 %0 = nvvm.match.sync any %thread_mask, %val32 : i32 -> !llvm.struct<(i32, i1)>172 llvm.return173}174 175// -----176 177llvm.func @nvvm_cvt_float_to_f8x2_invalid_rounding_e4m3(%a : f32, %b : f32) {178 // expected-error @below {{Only RN rounding mode is supported for conversions from f32x2 to 'f8E4M3FN' and 'f8E5M2' types}}179 %res = nvvm.convert.f32x2.to.f8x2 %a, %b {rnd = #nvvm.fp_rnd_mode<rz>, sat = #nvvm.sat_mode<satfinite>} : i16 (f8E4M3FN)180 llvm.return181}182 183// -----184 185llvm.func @nvvm_cvt_float_to_f8x2_invalid_rounding_e5m2(%a : f32, %b : f32) {186 // expected-error @below {{Only RN rounding mode is supported for conversions from f32x2 to 'f8E4M3FN' and 'f8E5M2' types}}187 %res = nvvm.convert.f32x2.to.f8x2 %a, %b {rnd = #nvvm.fp_rnd_mode<rp>, sat = #nvvm.sat_mode<satfinite>} : i16 (f8E5M2)188 llvm.return189}190 191// -----192 193llvm.func @nvvm_cvt_float_to_f8x2_invalid_rounding_ue8m0(%a : f32, %b : f32) {194 // expected-error @below {{Only RZ and RP rounding modes are supported for conversions from f32x2 to 'f8E8M0FNU' type}}195 %res = nvvm.convert.f32x2.to.f8x2 %a, %b {rnd = #nvvm.fp_rnd_mode<rn>} : i16 (f8E8M0FNU)196 llvm.return197}198 199// -----200 201llvm.func @nvvm_cvt_float_to_f8x2_invalid_saturation_e4m3(%a : f32, %b : f32) {202 // expected-error @below {{Only SATFINITE saturation mode is supported for conversions from f32x2 to 'f8E4M3FN' and 'f8E5M2' types}}203 %res = nvvm.convert.f32x2.to.f8x2 %a, %b {rnd = #nvvm.fp_rnd_mode<rn>, sat = #nvvm.sat_mode<none>} : i16 (f8E4M3FN)204 llvm.return205}206 207// -----208 209llvm.func @nvvm_cvt_float_to_f8x2_invalid_saturation_e5m2(%a : f32, %b : f32) {210 // expected-error @below {{Only SATFINITE saturation mode is supported for conversions from f32x2 to 'f8E4M3FN' and 'f8E5M2' types}}211 %res = nvvm.convert.f32x2.to.f8x2 %a, %b {rnd = #nvvm.fp_rnd_mode<rn>, sat = #nvvm.sat_mode<none>} : i16 (f8E5M2)212 llvm.return213}214 215// -----216 217llvm.func @nvvm_cvt_float_to_f8x2_relu_not_supported_ue8m0(%a : f32, %b : f32) {218 // expected-error @below {{relu not supported for conversions to 'f8E8M0FNU' type}}219 %res = nvvm.convert.f32x2.to.f8x2 %a, %b {rnd = #nvvm.fp_rnd_mode<rp>, relu = true} : i16 (f8E8M0FNU)220 llvm.return221}222 223// -----224 225llvm.func @nvvm_cvt_f16x2_to_f8x2_invalid_type(%src : vector<2xf16>) {226 // expected-error @below {{Only 'f8E4M3FN' and 'f8E5M2' types are supported for conversions from f16x2 to f8x2.}}227 %res = nvvm.convert.f16x2.to.f8x2 %src : vector<2xf16> -> i16 (f8E8M0FNU)228 llvm.return229}230 231// -----232 233llvm.func @nvvm_cvt_bf16x2_to_f8x2_invalid_type(%src : vector<2xbf16>) {234 // expected-error @below {{Only 'f8E8M0FNU' type is supported for conversions from bf16x2 to f8x2.}}235 %res = nvvm.convert.bf16x2.to.f8x2 %src {rnd = #nvvm.fp_rnd_mode<rz>} : vector<2xbf16> -> i16 (f8E4M3FN)236 llvm.return237}238 239// -----240 241llvm.func @nvvm_cvt_bf16x2_to_f8x2_invalid_rounding(%src : vector<2xbf16>) {242 // expected-error @below {{Only RZ and RP rounding modes are supported for conversions from bf16x2 to f8x2.}}243 %res = nvvm.convert.bf16x2.to.f8x2 %src {rnd = #nvvm.fp_rnd_mode<rn>} : vector<2xbf16> -> i16 (f8E8M0FNU)244 llvm.return245}246 247// -----248 249llvm.func @nvvm_cvt_f32x2_to_f6x2_invalid_type(%a : f32, %b : f32) {250 // expected-error @below {{Only 'f6E2M3FN' and 'f6E3M2FN' types are supported for conversions from f32x2 to f6x2.}}251 %res = nvvm.convert.f32x2.to.f6x2 %a, %b : i16 (f8E8M0FNU)252 llvm.return253}254 255// -----256 257llvm.func @nvvm_cvt_f32x2_to_f4x2_invalid_type(%a : f32, %b : f32) {258 // expected-error @below {{Only 'f4E2M1FN' type is supported for conversions from f32x2 to f4x2.}}259 %res = nvvm.convert.f32x2.to.f4x2 %a, %b : i8 (f8E4M3FN)260 llvm.return261}262 263// -----264 265llvm.func @nvvm_cvt_f8x2_to_f16x2_invalid_type(%src : vector<2xi8>) {266 // expected-error @below {{Only 'f8E4M3FN' and 'f8E5M2' types are supported for conversions from f8x2 to f16x2.}}267 %res = nvvm.convert.f8x2.to.f16x2 %src : vector<2xi8> (f8E4M3) -> vector<2xf16>268 llvm.return269}270 271// -----272 273llvm.func @nvvm_cvt_f8x2_to_bf16x2_invalid_type(%src : vector<2xi8>) {274 // expected-error @below {{Only 'f8E8M0FNU' type is supported for conversions from f8x2 to bf16x2.}}275 %res = nvvm.convert.f8x2.to.bf16x2 %src : vector<2xi8> (f8E4M3FN) -> vector<2xbf16>276 llvm.return277}278 279// -----280 281llvm.func @nvvm_cvt_f6x2_to_f16x2_invalid_type(%src : vector<2xi8>) {282 // expected-error @below {{Only 'f6E2M3FN' and 'f6E3M2FN' types are supported for conversions from f6x2 to f16x2.}}283 %res = nvvm.convert.f6x2.to.f16x2 %src : vector<2xi8> (f8E4M3FN) -> vector<2xf16>284 llvm.return285}286 287// -----288 289llvm.func @nvvm_cvt_f4x2_to_f16x2_invalid_type(%src : i8) {290 // expected-error @below {{Only 'f4E2M1FN' type is supported for conversions from f4x2 to f16x2.}}291 %res = nvvm.convert.f4x2.to.f16x2 %src : i8 (f6E2M3FN) -> vector<2xf16>292 llvm.return293}294 295// -----296 297llvm.func @nvvm_prefetch_L1_with_evict_priority(%global_ptr: !llvm.ptr<1>) {298 // expected-error @below {{cache eviction priority supported only for cache level L2}}299 nvvm.prefetch level = L1, evict_priority = evict_last, %global_ptr : !llvm.ptr<1>300 llvm.return301}302 303// -----304 305llvm.func @nvvm_prefetch_L2_with_evict_last_invalid_addr_space(%local_ptr: !llvm.ptr<5>) {306 // expected-error @below {{cache eviction priority requires a global pointer}}307 nvvm.prefetch level = L2, evict_priority = evict_last, %local_ptr : !llvm.ptr<5>308 llvm.return309}310 311// -----312 313llvm.func @nvvm_prefetch_L2_with_evict_normal_invalid_addr_space(%local_ptr: !llvm.ptr<5>) {314 // expected-error @below {{cache eviction priority requires a global pointer}}315 nvvm.prefetch level = L2, evict_priority = evict_normal, %local_ptr : !llvm.ptr<5>316 llvm.return317}318 319// -----320 321llvm.func @nvvm_prefetch_L2_with_invalid_evict_first(%global_ptr: !llvm.ptr<1>) {322 // expected-error @below {{unsupported cache eviction priority, only evict_last and evict_normal are supported}}323 nvvm.prefetch level = L2, evict_priority = evict_first, %global_ptr : !llvm.ptr<1>324 llvm.return325}326 327// -----328 329llvm.func @nvvm_prefetch_L2_with_invalid_evict_unchanged(%global_ptr: !llvm.ptr<1>) {330 // expected-error @below {{unsupported cache eviction priority, only evict_last and evict_normal are supported}}331 nvvm.prefetch level = L2, evict_priority = evict_unchanged, %global_ptr : !llvm.ptr<1>332 llvm.return333}334 335// -----336 337llvm.func @nvvm_prefetch_L2_with_invalid_no_allocate(%global_ptr: !llvm.ptr<1>) {338 // expected-error @below {{unsupported cache eviction priority, only evict_last and evict_normal are supported}}339 nvvm.prefetch level = L2, evict_priority = no_allocate, %global_ptr : !llvm.ptr<1>340 llvm.return341}342 343// -----344 345llvm.func @nvvm_prefetch_uniform_with_L2(%gen_ptr: !llvm.ptr) {346 // expected-error @below {{unsupported cache level, the only supported uniform cache level is L1}}347 nvvm.prefetch level = L2 uniform, %gen_ptr : !llvm.ptr348 llvm.return349}350 351// -----352 353llvm.func @nvvm_prefetch_uniform_with_invalid_addr_space(%global_ptr: !llvm.ptr<1>) {354 // expected-error @below {{prefetch to uniform cache requires a generic pointer}}355 nvvm.prefetch level = L1 uniform, %global_ptr : !llvm.ptr<1>356 llvm.return357}358 359// -----360 361llvm.func @nvvm_prefetch_both_tensormap_and_cache_level(%gen_ptr: !llvm.ptr) {362 // expected-error @below {{cannot specify both tensormap and cache level}}363 nvvm.prefetch level = L1, tensormap, %gen_ptr : !llvm.ptr364 llvm.return365}366 367// -----368 369llvm.func @nvvm_prefetch_tensormap_invalid_addr_space(%global_ptr: !llvm.ptr<1>) {370 // expected-error @below {{prefetch tensormap requires a generic or constant pointer}}371 nvvm.prefetch tensormap, %global_ptr : !llvm.ptr<1>372 llvm.return373}374 375// -----376 377llvm.func @nvvm_prefetch_tensormap_with_evict_priority(%gen_ptr: !llvm.ptr) {378 // expected-error @below {{prefetch tensormap does not support eviction priority}}379 nvvm.prefetch tensormap, evict_priority = evict_last, %gen_ptr : !llvm.ptr380 llvm.return381}382 383// -----384 385llvm.func @nvvm_prefetch_tensormap_in_param_space_non_generic(%const_ptr: !llvm.ptr<4>) {386 // expected-error @below {{in_param_space can only be specified for a generic pointer}}387 nvvm.prefetch tensormap in_param_space, %const_ptr : !llvm.ptr<4>388 llvm.return389}390 391// -----392 393llvm.func @nvvm_prefetch_cache_level_invalid_addr_space(%const_ptr: !llvm.ptr<4>) {394 // expected-error @below {{prefetch to cache level requires a generic, global, or local pointer}}395 nvvm.prefetch level = L1, %const_ptr : !llvm.ptr<4>396 llvm.return397}398 399// -----400 401llvm.func @nvvm_prefetch_predicate_without_tensormap(%gen_ptr: !llvm.ptr, %pred: i1) {402 // expected-error @below {{predicate supported only on prefetch tensormap}}403 nvvm.prefetch level = L1, %gen_ptr, predicate = %pred : !llvm.ptr, i1404 llvm.return405}406 407// -----408 409llvm.func @nvvm_prefetch_no_level_or_tensormap(%gen_ptr: !llvm.ptr) {410 // expected-error @below {{requires specification of either cache level or tensormap}}411 nvvm.prefetch %gen_ptr : !llvm.ptr412 llvm.return413}414 415// -----416 417llvm.func @st_matrix(%arg0: !llvm.ptr<3>, %r1: i32, %r2: i32, %r3: i32, %r4: i32) {418 // expected-error@+1 {{'nvvm.stmatrix' op expected num attribute to be 1, 2 or 4}}419 nvvm.stmatrix %arg0, %r1, %r2, %r3 {layout = #nvvm.mma_layout<row>, shape = #nvvm.ld_st_matrix_shape<m = 8, n = 8>, eltType = #nvvm.ld_st_matrix_elt_type<b16>} : !llvm.ptr<3>, i32, i32, i32420 llvm.return421}422 423// -----424 425llvm.func @st_matrix(%arg0: !llvm.ptr<3>, %r1: i32, %r2: i32, %r3: i32, %r4: i32) {426 // expected-error@+1 {{'nvvm.stmatrix' op expected shape to be 8x8 or 16x8}}427 nvvm.stmatrix %arg0, %r1 {layout = #nvvm.mma_layout<row>, shape = #nvvm.ld_st_matrix_shape<m = 16, n = 16>, eltType = #nvvm.ld_st_matrix_elt_type<b16>} : !llvm.ptr<3>, i32428 llvm.return429}430 431// -----432 433llvm.func @st_matrix(%arg0: !llvm.ptr<3>, %r1: i32, %r2: i32, %r3: i32, %r4: i32) {434 // expected-error@+1 {{'nvvm.stmatrix' op expected element type to be B16 for 8x8 matrix}}435 nvvm.stmatrix %arg0, %r1 {layout = #nvvm.mma_layout<row>, shape = #nvvm.ld_st_matrix_shape<m = 8, n = 8>, eltType = #nvvm.ld_st_matrix_elt_type<b8>} : !llvm.ptr<3>, i32436 llvm.return437}438// -----439 440llvm.func @st_matrix(%arg0: !llvm.ptr<3>, %r1: i32, %r2: i32, %r3: i32, %r4: i32) {441 // expected-error@+1 {{'nvvm.stmatrix' op expected element type to be B8 for 16x8 matrix}}442 nvvm.stmatrix %arg0, %r1 {layout = #nvvm.mma_layout<col>, shape = #nvvm.ld_st_matrix_shape<m = 16, n = 8>, eltType = #nvvm.ld_st_matrix_elt_type<b16>} : !llvm.ptr<3>, i32443 llvm.return444}445 446// -----447 448llvm.func @st_matrix(%arg0: !llvm.ptr<3>, %r1: i32, %r2: i32, %r3: i32, %r4: i32) {449 // expected-error@+1 {{'nvvm.stmatrix' op expected layout to be col for 16x8 matrix}}450 nvvm.stmatrix %arg0, %r1 {layout = #nvvm.mma_layout<row>, shape = #nvvm.ld_st_matrix_shape<m = 16, n = 8>, eltType = #nvvm.ld_st_matrix_elt_type<b8>} : !llvm.ptr<3>, i32451 llvm.return452}453 454// -----455 456llvm.func @st_matrix(%arg0: !llvm.ptr<3>, %r1: i32, %r2: i32, %r3: i32, %r4: i32) {457 // expected-error@+1 {{'nvvm.stmatrix' op expected num attribute to be 1, 2 or 4}}458 nvvm.stmatrix %arg0, %r1, %r2, %r3 {layout = #nvvm.mma_layout<row>, shape = #nvvm.ld_st_matrix_shape<m = 8, n = 8>, eltType = #nvvm.ld_st_matrix_elt_type<b16>} : !llvm.ptr<3>, i32, i32, i32459 llvm.return460}461 462// -----463 464llvm.func @st_matrix(%arg0: !llvm.ptr<3>, %r1: i32, %r2: i32, %r3: i32, %r4: i32) {465 // expected-error@+1 {{'nvvm.stmatrix' op expected shape to be 8x8 or 16x8}}466 nvvm.stmatrix %arg0, %r1 {layout = #nvvm.mma_layout<row>, shape = #nvvm.ld_st_matrix_shape<m = 16, n = 16>, eltType = #nvvm.ld_st_matrix_elt_type<b16>} : !llvm.ptr<3>, i32467 llvm.return468}469 470// -----471 472llvm.func @st_matrix(%arg0: !llvm.ptr<3>, %r1: i32, %r2: i32, %r3: i32, %r4: i32) {473 // expected-error@+1 {{'nvvm.stmatrix' op expected element type to be B16 for 8x8 matrix}}474 nvvm.stmatrix %arg0, %r1 {layout = #nvvm.mma_layout<row>, shape = #nvvm.ld_st_matrix_shape<m = 8, n = 8>, eltType = #nvvm.ld_st_matrix_elt_type<b8>} : !llvm.ptr<3>, i32475 llvm.return476}477// -----478 479llvm.func @st_matrix(%arg0: !llvm.ptr<3>, %r1: i32, %r2: i32, %r3: i32, %r4: i32) {480 // expected-error@+1 {{'nvvm.stmatrix' op expected element type to be B8 for 16x8 matrix}}481 nvvm.stmatrix %arg0, %r1 {layout = #nvvm.mma_layout<col>, shape = #nvvm.ld_st_matrix_shape<m = 16, n = 8>, eltType = #nvvm.ld_st_matrix_elt_type<b16>} : !llvm.ptr<3>, i32482 llvm.return483}484 485// -----486 487llvm.func @st_matrix(%arg0: !llvm.ptr<3>, %r1: i32, %r2: i32, %r3: i32, %r4: i32) {488 // expected-error@+1 {{'nvvm.stmatrix' op expected layout to be col for 16x8 matrix}}489 nvvm.stmatrix %arg0, %r1 {layout = #nvvm.mma_layout<row>, shape = #nvvm.ld_st_matrix_shape<m = 16, n = 8>, eltType = #nvvm.ld_st_matrix_elt_type<b8>} : !llvm.ptr<3>, i32490 llvm.return491}492 493// -----494 495llvm.func @ld_matrix(%arg0: !llvm.ptr<3>) {496 // expected-error@+1 {{'nvvm.ldmatrix' op expected num attribute to be 1, 2 or 4 for 8x8 matrix}}497 %l = nvvm.ldmatrix %arg0 {num = 3 : i32, layout = #nvvm.mma_layout<row>, shape = #nvvm.ld_st_matrix_shape<m = 8, n = 8>, eltType = #nvvm.ld_st_matrix_elt_type<b16>} : (!llvm.ptr<3>) -> i32498 llvm.return499}500 501// -----502 503llvm.func @ld_matrix(%arg0: !llvm.ptr<3>) {504 // expected-error@+1 {{'nvvm.ldmatrix' op expected destination type is i32}}505 %l = nvvm.ldmatrix %arg0 {num = 1 : i32, layout = #nvvm.mma_layout<row>, shape = #nvvm.ld_st_matrix_shape<m = 8, n = 8>, eltType = #nvvm.ld_st_matrix_elt_type<b16>} : (!llvm.ptr<3>) -> !llvm.struct<(i32)>506 llvm.return507}508 509// -----510 511llvm.func @ld_matrix(%arg0: !llvm.ptr<3>) {512 // expected-error@+1 {{'nvvm.ldmatrix' op expected destination type is a structure of 4 elements of type i32}}513 %l = nvvm.ldmatrix %arg0 {num = 4 : i32, layout = #nvvm.mma_layout<row>, shape = #nvvm.ld_st_matrix_shape<m = 8, n = 8>, eltType = #nvvm.ld_st_matrix_elt_type<b16>} : (!llvm.ptr<3>) -> !llvm.struct<(i32, i32)>514 llvm.return515}516 517// -----518 519llvm.func @ld_matrix(%arg0: !llvm.ptr<3>) {520 // expected-error@+1 {{'nvvm.ldmatrix' op expected element type to be b16 for 8x8 matrix}}521 %l = nvvm.ldmatrix %arg0 {num = 1 : i32, layout = #nvvm.mma_layout<row>, shape = #nvvm.ld_st_matrix_shape<m = 8, n = 8>, eltType = #nvvm.ld_st_matrix_elt_type<b8>} : (!llvm.ptr<3>) -> i32522 llvm.return523}524 525// -----526 527llvm.func @ld_matrix(%arg0: !llvm.ptr<3>) {528 // expected-error@+1 {{'nvvm.ldmatrix' op expected num attribute to be 1, 2 or 4 for 8x16 matrix}}529 %l = nvvm.ldmatrix %arg0 {num = 3 : i32, layout = #nvvm.mma_layout<row>, shape = #nvvm.ld_st_matrix_shape<m = 8, n = 16>, eltType = #nvvm.ld_st_matrix_elt_type<b8x16.b4x16_p64>} : (!llvm.ptr<3>) -> !llvm.struct<(i32, i32, i32)>530 llvm.return531}532 533// -----534 535llvm.func @ld_matrix(%arg0: !llvm.ptr<3>) {536 // expected-error@+1 {{'nvvm.ldmatrix' op expected layout to be row for 8x16 matrix}}537 %l = nvvm.ldmatrix %arg0 {num = 1 : i32, layout = #nvvm.mma_layout<col>, shape = #nvvm.ld_st_matrix_shape<m = 8, n = 16>, eltType = #nvvm.ld_st_matrix_elt_type<b8x16.b4x16_p64>} : (!llvm.ptr<3>) -> i32538 llvm.return539}540 541// -----542 543llvm.func @ld_matrix(%arg0: !llvm.ptr<3>) {544 // expected-error@+1 {{'nvvm.ldmatrix' op expected element type to be b8x16.b4x16_p64 or b8x16.b6x16_p32 for 8x16 matrix}}545 %l = nvvm.ldmatrix %arg0 {num = 1 : i32, layout = #nvvm.mma_layout<row>, shape = #nvvm.ld_st_matrix_shape<m = 8, n = 16>, eltType = #nvvm.ld_st_matrix_elt_type<b8>} : (!llvm.ptr<3>) -> i32546 llvm.return547}548 549// -----550 551llvm.func @ld_matrix(%arg0: !llvm.ptr<3>) {552 // expected-error@+1 {{'nvvm.ldmatrix' op expected num attribute to be 1 or 2 for 16x16 matrix}}553 %l = nvvm.ldmatrix %arg0 {num = 4 : i32, layout = #nvvm.mma_layout<row>, shape = #nvvm.ld_st_matrix_shape<m = 16, n = 16>, eltType = #nvvm.ld_st_matrix_elt_type<b8x16.b4x16_p64>} : (!llvm.ptr<3>) -> !llvm.struct<(i32, i32, i32, i32)>554 llvm.return555}556 557// -----558 559llvm.func @ld_matrix(%arg0: !llvm.ptr<3>) {560 // expected-error@+1 {{'nvvm.ldmatrix' op expected layout to be col for 16x16 matrix}}561 %l = nvvm.ldmatrix %arg0 {num = 1 : i32, layout = #nvvm.mma_layout<row>, shape = #nvvm.ld_st_matrix_shape<m = 16, n = 16>, eltType = #nvvm.ld_st_matrix_elt_type<b8x16.b6x16_p32>} : (!llvm.ptr<3>) -> !llvm.struct<(i32, i32, i32, i32)>562 llvm.return563}564 565// -----566 567llvm.func @ld_matrix(%arg0: !llvm.ptr<3>) {568 // expected-error@+1 {{'nvvm.ldmatrix' op expected element type to be b8, b8x16.b4x16_p64 or b8x16.b6x16_p32 for 16x16 matrix}}569 %l = nvvm.ldmatrix %arg0 {num = 1 : i32, layout = #nvvm.mma_layout<col>, shape = #nvvm.ld_st_matrix_shape<m = 16, n = 16>, eltType = #nvvm.ld_st_matrix_elt_type<b16>} : (!llvm.ptr<3>) -> i32570 llvm.return571}572 573llvm.func @ld_matrix(%arg0: !llvm.ptr<3>) {574 // expected-error@+1 {{'nvvm.ldmatrix' op expected destination type is a structure of 2 elements of type i32}}575 %l = nvvm.ldmatrix %arg0 {num = 1 : i32, layout = #nvvm.mma_layout<col>, shape = #nvvm.ld_st_matrix_shape<m = 16, n = 16>, eltType = #nvvm.ld_st_matrix_elt_type<b8>} : (!llvm.ptr<3>) -> i32576 llvm.return577}578 579// -----580 581llvm.func @clusterlaunchcontrol_query_cancel_is_canceled_invalid_return_type(%try_cancel_response: i128) {582 // expected-error@+1 {{'nvvm.clusterlaunchcontrol.query.cancel' op is_canceled query type returns an i1}}583 %res = nvvm.clusterlaunchcontrol.query.cancel query = is_canceled, %try_cancel_response : i32584 llvm.return585}586 587// -----588 589llvm.func @clusterlaunchcontrol_query_cancel_get_first_cta_id_invalid_return_type(%try_cancel_response: i128) {590 // expected-error@+1 {{'nvvm.clusterlaunchcontrol.query.cancel' op get_first_cta_id_x, get_first_cta_id_y, get_first_cta_id_z query types return an i32}}591 %res = nvvm.clusterlaunchcontrol.query.cancel query = get_first_cta_id_x, %try_cancel_response : i1592 llvm.return593}594 595// -----596 597// Test that ensures invalid row/col layouts for matrices A and B are not accepted598llvm.func @nvvm_mma_m16n8k32_s4_s4(%a0 : i32, %a1 : i32, %b0 : i32, %c0 : i32, %c1 : i32, %c2 : i32, %c3 : i32) -> !llvm.struct<(i32,i32,i32,i32)> {599 // expected-error@+1 {{Only m8n8k4 with f16 supports other layouts.}}600 %0 = nvvm.mma.sync A[%a0, %a1] B[%b0] C[%c0, %c1, %c2, %c3]601 {layoutA = #nvvm.mma_layout<col>, layoutB = #nvvm.mma_layout<col>,602 multiplicandAPtxType = #nvvm.mma_type<s4>, multiplicandBPtxType = #nvvm.mma_type<s4>,603 intOverflowBehavior=#nvvm.mma_int_overflow<satfinite>,604 shape = #nvvm.shape<m = 16, n = 8, k = 32>} : (i32, i32, i32) -> !llvm.struct<(i32,i32,i32,i32)>605 llvm.return %0 : !llvm.struct<(i32,i32,i32,i32)>606}607 608// -----609 610// Test for range validation - invalid range where lower == upper but not at extremes611func.func @invalid_range_equal_bounds() {612 // expected-error @below {{invalid range attribute: Lower == Upper, but they aren't min (0) or max (4294967295) value! This is an invalid constant range.}}613 %0 = nvvm.read.ptx.sreg.warpsize range <i32, 32, 32> : i32614 return615}616 617// -----618 619// Test for correct return type check for wmma.load fragment a for f64 620llvm.func @nvvm_wmma_load_a_f64(%arg0: !llvm.ptr, %arg1 : i32) {621 // expected-error @below {{'nvvm.wmma.load' op expected destination type to be f64}}622 %0 = nvvm.wmma.load %arg0, %arg1623 {eltype = #nvvm.mma_type<f64>, frag = #nvvm.mma_frag<a>, k = 4 : i32, layout = #nvvm.mma_layout<row>, m = 8 : i32, n = 8 : i32}624 : (!llvm.ptr) -> !llvm.struct<(f64)>625 llvm.return626}627