643 lines · plain
1// RUN: mlir-opt %s -split-input-file -verify-diagnostics | FileCheck %s2 3// CHECK-LABEL: @nvvm_special_regs4func.func @nvvm_special_regs() -> i32 {5 // CHECK: nvvm.read.ptx.sreg.tid.x : i326 %0 = nvvm.read.ptx.sreg.tid.x : i327 // CHECK: nvvm.read.ptx.sreg.tid.y : i328 %1 = nvvm.read.ptx.sreg.tid.y : i329 // CHECK: nvvm.read.ptx.sreg.tid.z : i3210 %2 = nvvm.read.ptx.sreg.tid.z : i3211 // CHECK: nvvm.read.ptx.sreg.ntid.x : i3212 %3 = nvvm.read.ptx.sreg.ntid.x : i3213 // CHECK: nvvm.read.ptx.sreg.ntid.y : i3214 %4 = nvvm.read.ptx.sreg.ntid.y : i3215 // CHECK: nvvm.read.ptx.sreg.ntid.z : i3216 %5 = nvvm.read.ptx.sreg.ntid.z : i3217 // CHECK: nvvm.read.ptx.sreg.ctaid.x : i3218 %6 = nvvm.read.ptx.sreg.ctaid.x : i3219 // CHECK: nvvm.read.ptx.sreg.ctaid.y : i3220 %7 = nvvm.read.ptx.sreg.ctaid.y : i3221 // CHECK: nvvm.read.ptx.sreg.ctaid.z : i3222 %8 = nvvm.read.ptx.sreg.ctaid.z : i3223 // CHECK: nvvm.read.ptx.sreg.nctaid.x : i3224 %9 = nvvm.read.ptx.sreg.nctaid.x : i3225 // CHECK: nvvm.read.ptx.sreg.nctaid.y : i3226 %10 = nvvm.read.ptx.sreg.nctaid.y : i3227 // CHECK: nvvm.read.ptx.sreg.nctaid.z : i3228 %11 = nvvm.read.ptx.sreg.nctaid.z : i3229 llvm.return %0 : i3230}31 32// CHECK-LABEL: @nvvm_rcp33func.func @nvvm_rcp(%arg0: f32) -> f32 {34 // CHECK: nvvm.rcp.approx.ftz.f %arg0 : f3235 %0 = nvvm.rcp.approx.ftz.f %arg0 : f3236 llvm.return %0 : f3237}38 39// CHECK-LABEL: @llvm_nvvm_barrier040func.func @llvm_nvvm_barrier0() {41 // CHECK: nvvm.barrier042 nvvm.barrier043 llvm.return44}45 46// CHECK-LABEL: @llvm_nvvm_barrier47// CHECK-SAME: (%[[barId:.*]]: i32, %[[numberOfThreads:.*]]: i32)48llvm.func @llvm_nvvm_barrier(%barId : i32, %numberOfThreads : i32) {49 // CHECK: nvvm.barrier 50 nvvm.barrier 51 // CHECK: nvvm.barrier id = %[[barId]]52 nvvm.barrier id = %barId53 // CHECK: nvvm.barrier id = %[[barId]] number_of_threads = %[[numberOfThreads]]54 nvvm.barrier id = %barId number_of_threads = %numberOfThreads55 llvm.return56}57 58// CHECK-LABEL: @llvm_nvvm_barrier_arrive59// CHECK-SAME: (%[[barId:.*]]: i32, %[[numberOfThreads:.*]]: i32)60llvm.func @llvm_nvvm_barrier_arrive(%barId : i32, %numberOfThreads : i32) {61 // CHECK: nvvm.barrier.arrive number_of_threads = %[[numberOfThreads]]62 nvvm.barrier.arrive number_of_threads = %numberOfThreads63 // CHECK: nvvm.barrier.arrive id = %[[barId]] number_of_threads = %[[numberOfThreads]]64 nvvm.barrier.arrive id = %barId number_of_threads = %numberOfThreads65 llvm.return66}67 68// CHECK-LABEL: @llvm_nvvm_cluster_arrive69func.func @llvm_nvvm_cluster_arrive() {70 // CHECK: nvvm.cluster.arrive71 nvvm.cluster.arrive72 // CHECK: nvvm.cluster.arrive {aligned}73 nvvm.cluster.arrive {aligned}74 llvm.return75}76 77// CHECK-LABEL: @llvm_nvvm_cluster_arrive_relaxed78func.func @llvm_nvvm_cluster_arrive_relaxed() {79 // CHECK: nvvm.cluster.arrive.relaxed80 nvvm.cluster.arrive.relaxed81 // CHECK: nvvm.cluster.arrive.relaxed {aligned}82 nvvm.cluster.arrive.relaxed {aligned}83 llvm.return84}85 86// CHECK-LABEL: @llvm_nvvm_cluster_wait87func.func @llvm_nvvm_cluster_wait() {88 // CHECK: nvvm.cluster.wait89 nvvm.cluster.wait90 // CHECK: nvvm.cluster.wait {aligned}91 nvvm.cluster.wait {aligned}92 llvm.return93}94 95// CHECK-LABEL: @llvm_nvvm_fence_sc_cluster96func.func @llvm_nvvm_fence_sc_cluster() {97 // CHECK: nvvm.fence.sc.cluster98 nvvm.fence.sc.cluster99 llvm.return100}101 102// CHECK-LABEL: @nvvm_shfl103func.func @nvvm_shfl(104 %arg0 : i32, %arg1 : i32, %arg2 : i32,105 %arg3 : i32, %arg4 : f32) -> i32 {106 // CHECK: nvvm.shfl.sync bfly %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}} : i32 -> i32107 %0 = nvvm.shfl.sync bfly %arg0, %arg3, %arg1, %arg2 : i32 -> i32108 // CHECK: nvvm.shfl.sync bfly %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}} : f32 -> f32109 %1 = nvvm.shfl.sync bfly %arg0, %arg4, %arg1, %arg2 : f32 -> f32110 // CHECK: nvvm.shfl.sync up %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}} : f32 -> f32111 %2 = nvvm.shfl.sync up %arg0, %arg4, %arg1, %arg2 : f32 -> f32112 // CHECK: nvvm.shfl.sync down %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}} : f32 -> f32113 %3 = nvvm.shfl.sync down %arg0, %arg4, %arg1, %arg2 : f32 -> f32114 // CHECK: nvvm.shfl.sync idx %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}} : f32 -> f32115 %4 = nvvm.shfl.sync idx %arg0, %arg4, %arg1, %arg2 : f32 -> f32116 llvm.return %0 : i32117}118 119// CHECK-LABEL: @nvvm_shfl_pred120func.func @nvvm_shfl_pred(121 %arg0 : i32, %arg1 : i32, %arg2 : i32,122 %arg3 : i32, %arg4 : f32) -> !llvm.struct<(i32, i1)> {123 // CHECK: nvvm.shfl.sync bfly %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}} {return_value_and_is_valid} : i32 -> !llvm.struct<(i32, i1)>124 %0 = nvvm.shfl.sync bfly %arg0, %arg3, %arg1, %arg2 {return_value_and_is_valid} : i32 -> !llvm.struct<(i32, i1)>125 // CHECK: nvvm.shfl.sync bfly %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}} {return_value_and_is_valid} : f32 -> !llvm.struct<(f32, i1)>126 %1 = nvvm.shfl.sync bfly %arg0, %arg4, %arg1, %arg2 {return_value_and_is_valid} : f32 -> !llvm.struct<(f32, i1)>127 llvm.return %0 : !llvm.struct<(i32, i1)>128}129 130// CHECK-LABEL: @nvvm_vote(131func.func @nvvm_vote(%arg0 : i32, %arg1 : i1) -> i32 {132 // CHECK: nvvm.vote.sync ballot %{{.*}}, %{{.*}} -> i32133 %0 = nvvm.vote.sync ballot %arg0, %arg1 -> i32134 // CHECK: nvvm.vote.sync all %{{.*}}, %{{.*}} -> i1135 %1 = nvvm.vote.sync all %arg0, %arg1 -> i1136 // CHECK: nvvm.vote.sync any %{{.*}}, %{{.*}} -> i1137 %2 = nvvm.vote.sync any %arg0, %arg1 -> i1138 // CHECK: nvvm.vote.sync uni %{{.*}}, %{{.*}} -> i1139 %3 = nvvm.vote.sync uni %arg0, %arg1 -> i1140 llvm.return %0 : i32141}142 143// CHECK-LABEL: @llvm_nvvm_bar_warp_sync144func.func @llvm_nvvm_bar_warp_sync(%mask : i32) {145 // CHECK: nvvm.bar.warp.sync %{{.*}}146 nvvm.bar.warp.sync %mask : i32147 llvm.return148}149 150// CHECK-LABEL: @nvvm_mma_m8n8k4_row_col_f32_f32151func.func @nvvm_mma_m8n8k4_row_col_f32_f32(%a0 : vector<2xf16>, %a1 : vector<2xf16>,152 %b0 : vector<2xf16>, %b1 : vector<2xf16>,153 %c0 : f32, %c1 : f32, %c2 : f32, %c3 : f32, %c4 : f32, %c5 : f32, %c6 : f32, %c7 : f32) {154 // CHECK: nvvm.mma.sync155 %0 = nvvm.mma.sync A[%a0, %a1] B[%b0, %b1] C[%c0, %c1, %c2, %c3, %c4, %c5, %c6, %c7]156 {layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>,157 shape = #nvvm.shape<m = 8, n = 8, k = 4>} : (vector<2xf16>, vector<2xf16>, f32) -> !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)>158 llvm.return %0 : !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)>159}160 161// CHECK-LABEL: @nvvm_mma_m8n8k4_f16_f16162func.func @nvvm_mma_m8n8k4_f16_f16(%a0 : vector<2xf16>, %a1 : vector<2xf16>,163 %b0 : vector<2xf16>, %b1 : vector<2xf16>,164 %c0 : vector<2xf16>, %c1 : vector<2xf16>, %c2 : vector<2xf16>, %c3 : vector<2xf16>) {165 // CHECK: nvvm.mma.sync A[{{.*}}] B[{{.*}}] C[{{.*}}]166 %0 = nvvm.mma.sync A[%a0, %a1] B[%b0, %b1] C[%c0, %c1, %c2, %c3]167 {layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>,168 shape = #nvvm.shape<m = 8, n = 8, k = 4>} : (vector<2xf16>,vector<2xf16>,vector<2xf16>) -> !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>169 llvm.return %0 : !llvm.struct<(vector<2xf16>, vector<2xf16>, vector<2xf16>, vector<2xf16>)>170}171 172// CHECK-LABEL: @nvvm_mma_m16n8k8_bf16_bf16173func.func @nvvm_mma_m16n8k8_bf16_bf16(%a0 : i32, %a1 : i32, %b0 : i32,174 %c0 : f32, %c1 : f32, %c2 : f32, %c3 : f32) {175 // CHECK: nvvm.mma.sync A[{{.*}}] B[{{.*}}] C[{{.*}}] {layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>, multiplicandAPtxType = #nvvm.mma_type<bf16>, multiplicandBPtxType = #nvvm.mma_type<bf16>, shape = #nvvm.shape<m = 16, n = 8, k = 8>} : (i32, i32, f32) -> !llvm.struct<(f32, f32, f32, f32)>176 %0 = nvvm.mma.sync A[%a0, %a1] B[%b0] C[%c0, %c1, %c2, %c3]177 {layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>,178 multiplicandAPtxType = #nvvm.mma_type<bf16>, multiplicandBPtxType = #nvvm.mma_type<bf16>,179 shape = #nvvm.shape<m = 16, n = 8, k = 8>} : (i32, i32, f32) -> !llvm.struct<(f32, f32, f32, f32)>180 llvm.return %0 : !llvm.struct<(f32, f32, f32, f32)>181}182 183// CHECK-LABEL: @nvvm_mma_m16n8k16_bf16_bf16184func.func @nvvm_mma_m16n8k16_bf16_bf16(%a0 : i32, %a1 : i32, %a2 : i32, %a3 : i32,185 %b0 : i32, %b1 : i32,186 %c0 : f32, %c1 : f32, %c2 : f32, %c3 : f32) {187 // CHECK: nvvm.mma.sync A[{{.*}}] B[{{.*}}] C[{{.*}}] {layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>, multiplicandAPtxType = #nvvm.mma_type<bf16>, multiplicandBPtxType = #nvvm.mma_type<bf16>, shape = #nvvm.shape<m = 16, n = 8, k = 16>} : (i32, i32, f32) -> !llvm.struct<(f32, f32, f32, f32)>188 %0 = nvvm.mma.sync A[%a0, %a1, %a2, %a3] B[%b0, %b1] C[%c0, %c1, %c2, %c3]189 {layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>,190 multiplicandAPtxType = #nvvm.mma_type<bf16>, multiplicandBPtxType = #nvvm.mma_type<bf16>,191 shape = #nvvm.shape<m = 16, n = 8, k = 16>} : (i32, i32, f32) -> !llvm.struct<(f32, f32, f32, f32)>192 llvm.return %0 : !llvm.struct<(f32, f32, f32, f32)>193}194 195// CHECK-LABEL: @nvvm_mma_m8n8k16_s8_s8196func.func @nvvm_mma_m8n8k16_s8_s8(%a0 : i32, %b0 : i32,197 %c0 : i32, %c1 : i32) {198 // CHECK: nvvm.mma.sync A[{{.*}}] B[{{.*}}] C[{{.*}}, {{.*}}] {intOverflowBehavior = #nvvm.mma_int_overflow<wrapped>, layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>, multiplicandAPtxType = #nvvm.mma_type<s8>, multiplicandBPtxType = #nvvm.mma_type<s8>, shape = #nvvm.shape<m = 8, n = 8, k = 16>} : (i32, i32, i32) -> !llvm.struct<(i32, i32)>199 %0 = nvvm.mma.sync A[%a0] B[%b0] C[%c0, %c1]200 {layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>,201 multiplicandAPtxType = #nvvm.mma_type<s8>, multiplicandBPtxType = #nvvm.mma_type<s8>,202 intOverflowBehavior=#nvvm.mma_int_overflow<wrapped>,203 shape = #nvvm.shape<m = 8, n = 8, k = 16>} : (i32, i32, i32) -> !llvm.struct<(i32, i32)>204 llvm.return %0 : !llvm.struct<(i32, i32)>205}206 207// CHECK-LABEL: @nvvm_mma_m16n8k8_f16_f16208func.func @nvvm_mma_m16n8k8_f16_f16(%a0 : vector<2xf16>, %a1 : vector<2xf16>,209 %b0 : vector<2xf16>,210 %c0 : vector<2xf16>, %c1 : vector<2xf16>) {211 // CHECK: nvvm.mma.sync A[%{{.*}}, %{{.*}}] B[%{{.*}}] C[%{{.*}}, %{{.*}}] {{{.*}}} : (vector<2xf16>, vector<2xf16>, vector<2xf16>) -> !llvm.struct<(vector<2xf16>, vector<2xf16>)>212 %0 = nvvm.mma.sync A[%a0, %a1] B[%b0] C[%c0, %c1]213 {layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>,214 shape = #nvvm.shape<m = 16, n = 8, k = 8>} : (vector<2xf16>, vector<2xf16>, vector<2xf16>) -> !llvm.struct<(vector<2xf16>, vector<2xf16>)>215 llvm.return %0 : !llvm.struct<(vector<2xf16>, vector<2xf16>)>216}217 218// CHECK-LABEL: @nvvm_mma_m16n8k16_f16_f16219func.func @nvvm_mma_m16n8k16_f16_f16(%a0 : vector<2xf16>, %a1 : vector<2xf16>,220 %a2 : vector<2xf16>, %a3 : vector<2xf16>,221 %b0 : vector<2xf16>, %b1 : vector<2xf16>,222 %c0 : vector<2xf16>, %c1 : vector<2xf16>) {223 // CHECK: nvvm.mma.sync A[{{.*}}, {{.*}}, {{.*}}, {{.*}}] B[{{.*}}, {{.*}}] C[{{.*}}, {{.*}}] {layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>, shape = #nvvm.shape<m = 16, n = 8, k = 16>} : (vector<2xf16>, vector<2xf16>, vector<2xf16>) -> !llvm.struct<(vector<2xf16>, vector<2xf16>)>224 %0 = nvvm.mma.sync A[%a0, %a1, %a2, %a3] B[%b0, %b1] C[%c0, %c1]225 {layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>,226 shape = #nvvm.shape<m = 16, n = 8, k = 16>} : (vector<2xf16>, vector<2xf16>, vector<2xf16>) -> !llvm.struct<(vector<2xf16>, vector<2xf16>)>227 llvm.return %0 : !llvm.struct<(vector<2xf16>, vector<2xf16>)>228}229 230// CHECK-LABEL: @nvvm_mma_m16n8k16_f32_f32231func.func @nvvm_mma_m16n8k16_f32_f32(%a0 : vector<2xf16>, %a1 : vector<2xf16>,232 %a2 : vector<2xf16>, %a3 : vector<2xf16>,233 %b0 : vector<2xf16>, %b1 : vector<2xf16>,234 %c0 : f32, %c1 : f32, %c2 : f32, %c3 : f32) {235 // CHECK: nvvm.mma.sync A[{{.*}}, {{.*}}, {{.*}}, {{.*}}] B[{{.*}}, {{.*}}] C[{{.*}}, {{.*}}, {{.*}}, {{.*}}] {layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>, shape = #nvvm.shape<m = 16, n = 8, k = 16>} : (vector<2xf16>, vector<2xf16>, f32) -> !llvm.struct<(f32, f32, f32, f32)>236 %0 = nvvm.mma.sync A[%a0, %a1, %a2, %a3] B[%b0, %b1] C[%c0, %c1, %c2, %c3]237 {layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>,238 shape = #nvvm.shape<m = 16, n = 8, k = 16>} : (vector<2xf16>, vector<2xf16>, f32) -> !llvm.struct<(f32, f32, f32, f32)>239 llvm.return %0 : !llvm.struct<(f32, f32, f32, f32)>240}241 242// CHECK-LABEL: @nvvm_mma_m16n8k4_tf32_f32243func.func @nvvm_mma_m16n8k4_tf32_f32(%a0 : i32, %a1 : i32,244 %b0 : i32,245 %c0 : f32, %c1 : f32, %c2 : f32, %c3 : f32) {246 // CHECK: nvvm.mma.sync A[{{.*}}, {{.*}}] B[{{.*}}] C[{{.*}}, {{.*}}, {{.*}}, {{.*}}] {layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>, multiplicandAPtxType = #nvvm.mma_type<tf32>, multiplicandBPtxType = #nvvm.mma_type<tf32>, shape = #nvvm.shape<m = 16, n = 8, k = 4>} : (i32, i32, f32) -> !llvm.struct<(f32, f32, f32, f32)>247 %0 = nvvm.mma.sync A[%a0, %a1] B[%b0] C[%c0, %c1, %c2, %c3]248 {layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>,249 multiplicandAPtxType = #nvvm.mma_type<tf32>, multiplicandBPtxType = #nvvm.mma_type<tf32>,250 shape = #nvvm.shape<m = 16, n = 8, k = 4>} : (i32, i32, f32) -> !llvm.struct<(f32, f32, f32, f32)>251 llvm.return %0 : !llvm.struct<(f32, f32, f32, f32)>252}253 254// CHECK-LABEL: @nvvm_mma_m16n8k16_s8_s8255func.func @nvvm_mma_m16n8k16_s8_s8(%a0 : i32, %a1 : i32, %b0 : i32,256 %c0 : i32, %c1 : i32, %c2 : i32, %c3 : i32) {257 // CHECK: nvvm.mma.sync A[{{.*}}, {{.*}}] B[{{.*}}] C[{{.*}}, {{.*}}, {{.*}}, {{.*}}] {intOverflowBehavior = #nvvm.mma_int_overflow<wrapped>, layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>, multiplicandAPtxType = #nvvm.mma_type<s8>, multiplicandBPtxType = #nvvm.mma_type<s8>, shape = #nvvm.shape<m = 16, n = 8, k = 16>} : (i32, i32, i32) -> !llvm.struct<(i32, i32, i32, i32)>258 %0 = nvvm.mma.sync A[%a0, %a1] B[%b0] C[%c0, %c1, %c2, %c3]259 {layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>,260 multiplicandAPtxType = #nvvm.mma_type<s8>, multiplicandBPtxType = #nvvm.mma_type<s8>,261 intOverflowBehavior=#nvvm.mma_int_overflow<wrapped>,262 shape = #nvvm.shape<m = 16, n = 8, k = 16>} : (i32, i32, i32) -> !llvm.struct<(i32,i32,i32,i32)>263 llvm.return %0 : !llvm.struct<(i32,i32,i32,i32)>264}265 266// CHECK-LABEL: @nvvm_mma_m16n8k16_s8_u8267func.func @nvvm_mma_m16n8k16_s8_u8(%a0 : i32, %a1 : i32,268 %b0 : i32,269 %c0 : i32, %c1 : i32, %c2 : i32, %c3 : i32) {270 // CHECK: nvvm.mma.sync A[{{.*}}, {{.*}}] B[{{.*}}] C[{{.*}}, {{.*}}, {{.*}}, {{.*}}] {intOverflowBehavior = #nvvm.mma_int_overflow<satfinite>, layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>, multiplicandAPtxType = #nvvm.mma_type<s8>, multiplicandBPtxType = #nvvm.mma_type<u8>, shape = #nvvm.shape<m = 16, n = 8, k = 16>} : (i32, i32, i32) -> !llvm.struct<(i32, i32, i32, i32)>271 %0 = nvvm.mma.sync A[%a0, %a1] B[%b0] C[%c0, %c1, %c2, %c3]272 {layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>,273 multiplicandAPtxType = #nvvm.mma_type<s8>, multiplicandBPtxType = #nvvm.mma_type<u8>,274 intOverflowBehavior=#nvvm.mma_int_overflow<satfinite>,275 shape = #nvvm.shape<m = 16, n = 8, k = 16>} : (i32, i32, i32) -> !llvm.struct<(i32,i32,i32,i32)>276 llvm.return %0 : !llvm.struct<(i32,i32,i32,i32)>277}278 279// CHECK-LABEL: @nvvm_mma_m16n8k256_b1_b1280func.func @nvvm_mma_m16n8k256_b1_b1(%a0 : i32, %a1 : i32, %a2 : i32, %a3 : i32,281 %b0 : i32, %b1 : i32,282 %c0 : i32, %c1 : i32, %c2 : i32, %c3 : i32) {283 // CHECK: nvvm.mma.sync A[{{.*}}, {{.*}}, {{.*}}, {{.*}}] B[{{.*}}] C[{{.*}}, {{.*}}, {{.*}}, {{.*}}] {b1Op = #nvvm.mma_b1op<xor_popc>, layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>, multiplicandAPtxType = #nvvm.mma_type<b1>, multiplicandBPtxType = #nvvm.mma_type<b1>, shape = #nvvm.shape<m = 16, n = 8, k = 256>} : (i32, i32, i32) -> !llvm.struct<(i32, i32, i32, i32)>284 %0 = nvvm.mma.sync A[%a0, %a1, %a2, %a3] B[%b0, %b1] C[%c0, %c1, %c2, %c3]285 {layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>,286 multiplicandAPtxType = #nvvm.mma_type<b1>, multiplicandBPtxType = #nvvm.mma_type<b1>,287 b1Op = #nvvm.mma_b1op<xor_popc>, shape = #nvvm.shape<m = 16, n = 8, k = 256>} : (i32, i32, i32) -> !llvm.struct<(i32,i32,i32,i32)>288 llvm.return %0 : !llvm.struct<(i32,i32,i32,i32)>289}290 291// CHECK-LABEL: @nvvm_mma_m16n8k128_b1_b1292func.func @nvvm_mma_m16n8k128_b1_b1(%a0 : i32, %a1 : i32,293 %b0 : i32,294 %c0 : i32, %c1 : i32, %c2 : i32, %c3 : i32) {295 // CHECK: nvvm.mma.sync A[{{.*}}, {{.*}}] B[{{.*}}] C[{{.*}}, {{.*}}, {{.*}}, {{.*}}] {b1Op = #nvvm.mma_b1op<xor_popc>, layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>, multiplicandAPtxType = #nvvm.mma_type<b1>, multiplicandBPtxType = #nvvm.mma_type<b1>, shape = #nvvm.shape<m = 16, n = 8, k = 128>} : (i32, i32, i32) -> !llvm.struct<(i32, i32, i32, i32)>296 %0 = nvvm.mma.sync A[%a0, %a1] B[%b0] C[%c0, %c1, %c2, %c3]297 {layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>,298 multiplicandAPtxType = #nvvm.mma_type<b1>, multiplicandBPtxType = #nvvm.mma_type<b1>,299 b1Op = #nvvm.mma_b1op<xor_popc>,300 shape = #nvvm.shape<m = 16, n = 8, k = 128>} : (i32, i32, i32) -> !llvm.struct<(i32,i32,i32,i32)>301 llvm.return %0 : !llvm.struct<(i32,i32,i32,i32)>302}303 304// CHECK-LABEL: @nvvm_mma_m8n8k128_b1_b1305func.func @nvvm_mma_m8n8k128_b1_b1(%a0 : i32,306 %b0 : i32,307 %c0 : i32, %c1 : i32) {308 // CHECK: nvvm.mma.sync A[{{.*}}] B[{{.*}}] C[{{.*}}, {{.*}}] {b1Op = #nvvm.mma_b1op<xor_popc>, layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>, multiplicandAPtxType = #nvvm.mma_type<b1>, multiplicandBPtxType = #nvvm.mma_type<b1>, shape = #nvvm.shape<m = 8, n = 8, k = 128>} : (i32, i32, i32) -> !llvm.struct<(i32, i32)>309 %0 = nvvm.mma.sync A[%a0] B[%b0] C[%c0, %c1]310 {layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>,311 multiplicandAPtxType = #nvvm.mma_type<b1>, multiplicandBPtxType = #nvvm.mma_type<b1>,312 b1Op = #nvvm.mma_b1op<xor_popc>, shape = #nvvm.shape<m = 8, n = 8, k = 128>} : (i32, i32, i32) -> !llvm.struct<(i32,i32)>313 llvm.return %0 : !llvm.struct<(i32,i32)>314}315 316// CHECK-LABEL: @nvvm_mma_m16n8k32_s4_s4317func.func @nvvm_mma_m16n8k32_s4_s4(%a0 : i32, %a1 : i32,318 %b0 : i32,319 %c0 : i32, %c1 : i32, %c2 : i32, %c3 : i32) {320 // CHECK: nvvm.mma.sync A[{{.*}}, {{.*}}] B[{{.*}}] C[{{.*}}, {{.*}}, {{.*}}, {{.*}}] {intOverflowBehavior = #nvvm.mma_int_overflow<wrapped>, layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>, multiplicandAPtxType = #nvvm.mma_type<s4>, multiplicandBPtxType = #nvvm.mma_type<s4>, shape = #nvvm.shape<m = 16, n = 8, k = 32>} : (i32, i32, i32) -> !llvm.struct<(i32, i32, i32, i32)>321 %0 = nvvm.mma.sync A[%a0, %a1] B[%b0] C[%c0, %c1, %c2, %c3]322 {layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>,323 multiplicandAPtxType = #nvvm.mma_type<s4>, multiplicandBPtxType = #nvvm.mma_type<s4>,324 intOverflowBehavior=#nvvm.mma_int_overflow<wrapped>,325 shape = #nvvm.shape<m = 16, n = 8, k = 32>} : (i32, i32, i32) -> !llvm.struct<(i32,i32,i32,i32)>326 llvm.return %0 : !llvm.struct<(i32,i32,i32,i32)>327}328 329// CHECK-LABEL: @nvvm_wmma_load_tf32330func.func @nvvm_wmma_load_tf32(%arg0: !llvm.ptr, %arg1 : i32) -> !llvm.struct<(i32, i32, i32, i32)> {331 // CHECK: nvvm.wmma.load {{.*}} {eltype = #nvvm.mma_type<tf32>, frag = #nvvm.mma_frag<a>, k = 8 : i32, layout = #nvvm.mma_layout<row>, m = 16 : i32, n = 16 : i32}332 %0 = nvvm.wmma.load %arg0, %arg1333 {eltype = #nvvm.mma_type<tf32>, frag = #nvvm.mma_frag<a>, k = 8 : i32, layout = #nvvm.mma_layout<row>, m = 16 : i32, n = 16 : i32}334 : (!llvm.ptr) -> !llvm.struct<(i32, i32, i32, i32)>335 llvm.return %0 : !llvm.struct<(i32, i32, i32, i32)>336}337 338// CHECK-LABEL: @nvvm_wmma_mma339func.func @nvvm_wmma_mma(%0 : i32, %1 : i32, %2 : i32, %3 : i32, %4 : i32, %5 : i32,340 %6 : i32, %7 : i32, %8 : f32, %9 : f32, %10 : f32,341 %11 : f32, %12 : f32, %13 : f32, %14 : f32, %15 : f32)342 -> !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)> {343 // CHECK: nvvm.wmma.mma {{.*}} {eltypeA = #nvvm.mma_type<tf32>, eltypeB = #nvvm.mma_type<f32>, k = 8 : i32, layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<row>, m = 16 : i32, n = 16 : i32}344 %r = nvvm.wmma.mma %0, %1, %2, %3, %4, %5, %6, %7, %8, %9, %10, %11, %12, %13, %14, %15345 {eltypeA = #nvvm.mma_type<tf32>, eltypeB = #nvvm.mma_type<f32>, k = 8 : i32, layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<row>, m = 16 : i32, n = 16 : i32}346 : (i32, i32, i32, i32, i32, i32, i32, i32, f32, f32, f32, f32, f32, f32, f32, f32)347 -> !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)>348 llvm.return %r : !llvm.struct<(f32, f32, f32, f32, f32, f32, f32, f32)>349}350 351// CHECK-LABEL: @cp_async352llvm.func @cp_async(%arg0: !llvm.ptr<3>, %arg1: !llvm.ptr<1>) {353// CHECK: nvvm.cp.async.shared.global %{{.*}}, %{{.*}}, 16, cache = ca354 nvvm.cp.async.shared.global %arg0, %arg1, 16, cache = ca : !llvm.ptr<3>, !llvm.ptr<1>355// CHECK: nvvm.cp.async.shared.global %{{.*}}, %{{.*}}, 16, cache = cg356 nvvm.cp.async.shared.global %arg0, %arg1, 16, cache = cg : !llvm.ptr<3>, !llvm.ptr<1>357// CHECK: nvvm.cp.async.commit.group358 nvvm.cp.async.commit.group359// CHECK: nvvm.cp.async.wait.group 0360 nvvm.cp.async.wait.group 0361 llvm.return362}363 364// CHECK-LABEL: llvm.func @redux_sync365llvm.func @redux_sync(%value : i32, %offset : i32) -> i32 {366 // CHECK: nvvm.redux.sync add %{{.*}}367 %r1 = nvvm.redux.sync add %value, %offset : i32 -> i32368 // CHECK: nvvm.redux.sync max %{{.*}}369 %r2 = nvvm.redux.sync max %value, %offset : i32 -> i32370 // CHECK: nvvm.redux.sync min %{{.*}}371 %r3 = nvvm.redux.sync min %value, %offset : i32 -> i32372 // CHECK: nvvm.redux.sync umax %{{.*}}373 %r5 = nvvm.redux.sync umax %value, %offset : i32 -> i32374 // CHECK: nvvm.redux.sync umin %{{.*}}375 %r6 = nvvm.redux.sync umin %value, %offset : i32 -> i32376 // CHECK: nvvm.redux.sync and %{{.*}}377 %r7 = nvvm.redux.sync and %value, %offset : i32 -> i32378 // CHECK: nvvm.redux.sync or %{{.*}}379 %r8 = nvvm.redux.sync or %value, %offset : i32 -> i32380 // CHECK: nvvm.redux.sync xor %{{.*}}381 %r9 = nvvm.redux.sync xor %value, %offset : i32 -> i32382 llvm.return %r1 : i32383}384 385llvm.func @redux_sync_f32(%value: f32, %offset: i32) -> f32 {386 // CHECK: nvvm.redux.sync fmin %{{.*}}387 %r1 = nvvm.redux.sync fmin %value, %offset: f32 -> f32388 // CHECK: nvvm.redux.sync fmin %{{.*}}389 %r2 = nvvm.redux.sync fmin %value, %offset {abs = true}: f32 -> f32390 // CHECK: nvvm.redux.sync fmin %{{.*}}391 %r3 = nvvm.redux.sync fmin %value, %offset {NaN = true}: f32 -> f32392 // CHECK: nvvm.redux.sync fmin %{{.*}}393 %r4 = nvvm.redux.sync fmin %value, %offset {abs = true, NaN = true}: f32 -> f32394 // CHECK: nvvm.redux.sync fmax %{{.*}}395 %r5 = nvvm.redux.sync fmax %value, %offset: f32 -> f32396 // CHECK: nvvm.redux.sync fmax %{{.*}}397 %r6 = nvvm.redux.sync fmax %value, %offset {abs = true}: f32 -> f32398 // CHECK: nvvm.redux.sync fmax %{{.*}}399 %r7 = nvvm.redux.sync fmax %value, %offset {NaN = true}: f32 -> f32400 // CHECK: nvvm.redux.sync fmax %{{.*}}401 %r8 = nvvm.redux.sync fmax %value, %offset {abs = true, NaN = true}: f32 -> f32402 llvm.return %r1 : f32403}404 405// -----406 407// expected-error@below {{attribute attached to unexpected op}}408func.func private @expected_llvm_func() attributes { nvvm.kernel }409 410// -----411 412llvm.func private @mbarrier_init_generic(%barrier: !llvm.ptr) {413 %count = nvvm.read.ptx.sreg.ntid.x : i32414 // CHECK: nvvm.mbarrier.init %{{.*}}, %{{.*}} : !llvm.ptr, i32415 nvvm.mbarrier.init %barrier, %count : !llvm.ptr, i32416 llvm.return417}418 419 420llvm.func private @mbarrier_init_shared(%barrier: !llvm.ptr<3>) {421 %count = nvvm.read.ptx.sreg.ntid.x : i32422 // CHECK: nvvm.mbarrier.init %{{.*}}, %{{.*}} : !llvm.ptr<3>, i32423 nvvm.mbarrier.init %barrier, %count : !llvm.ptr<3>, i32424 llvm.return425}426 427 428llvm.func private @mbarrier_inval_generic(%barrier: !llvm.ptr) {429 // CHECK: nvvm.mbarrier.inval %{{.*}} : !llvm.ptr430 nvvm.mbarrier.inval %barrier : !llvm.ptr431 llvm.return432}433 434 435llvm.func private @mbarrier_inval_shared(%barrier: !llvm.ptr<3>) {436 // CHECK: nvvm.mbarrier.inval %{{.*}} : !llvm.ptr<3>437 nvvm.mbarrier.inval %barrier : !llvm.ptr<3>438 llvm.return439}440 441llvm.func private @mbarrier_arrive(%barrier: !llvm.ptr) {442 // CHECK: nvvm.mbarrier.arrive %{{.*}} : !llvm.ptr443 %0 = nvvm.mbarrier.arrive %barrier : !llvm.ptr -> i64444 llvm.return445}446 447llvm.func private @mbarrier_arrive_shared(%barrier: !llvm.ptr<3>) {448 // CHECK: nvvm.mbarrier.arrive %{{.*}} : !llvm.ptr<3>449 %0 = nvvm.mbarrier.arrive %barrier : !llvm.ptr<3> -> i64450 llvm.return451}452 453llvm.func private @mbarrier_arrive_nocomplete(%barrier: !llvm.ptr) {454 %count = nvvm.read.ptx.sreg.ntid.x : i32455 // CHECK: nvvm.mbarrier.arrive.nocomplete %{{.*}} : !llvm.ptr456 %0 = nvvm.mbarrier.arrive.nocomplete %barrier, %count : !llvm.ptr, i32 -> i64457 llvm.return458}459 460llvm.func private @mbarrier_arrive_nocomplete_shared(%barrier: !llvm.ptr<3>) {461 %count = nvvm.read.ptx.sreg.ntid.x : i32462 // CHECK: nvvm.mbarrier.arrive.nocomplete %{{.*}} : !llvm.ptr<3>463 %0 = nvvm.mbarrier.arrive.nocomplete %barrier, %count : !llvm.ptr<3>, i32 -> i64464 llvm.return465}466 467llvm.func private @mbarrier_test_wait(%barrier: !llvm.ptr, %token : i64) -> i1 { 468 // CHECK: nvvm.mbarrier.test.wait %{{.*}}469 %isComplete = nvvm.mbarrier.test.wait %barrier, %token : !llvm.ptr, i64 -> i1470 llvm.return %isComplete : i1471}472 473llvm.func private @mbarrier_test_wait_shared(%barrier: !llvm.ptr<3>, %token : i64) {474 %count = nvvm.read.ptx.sreg.ntid.x : i32475 // CHECK: nvvm.mbarrier.test.wait %{{.*}}476 %isComplete = nvvm.mbarrier.test.wait %barrier, %token : !llvm.ptr<3>, i64 -> i1477 llvm.return478}479 480// CHECK-LABEL: @wgmma_fence_aligned481func.func @wgmma_fence_aligned() {482 // CHECK: nvvm.wgmma.fence.aligned483 nvvm.wgmma.fence.aligned484 return485}486 487// CHECK-LABEL: @wgmma_commit_group_sync_aligned488func.func @wgmma_commit_group_sync_aligned() {489 // CHECK: nvvm.wgmma.commit.group.sync.aligned490 nvvm.wgmma.commit.group.sync.aligned491 return492}493 494 495// CHECK-LABEL: @wgmma_wait_group_sync_aligned496func.func @wgmma_wait_group_sync_aligned() {497 // CHECK: nvvm.wgmma.wait.group.sync.aligned498 nvvm.wgmma.wait.group.sync.aligned 0499 return500}501 502func.func @griddepcontrol_wait() {503 // CHECK: nvvm.griddepcontrol wait504 nvvm.griddepcontrol wait505 return506}507 508func.func @griddepcontrol_launch_dependents()509{510 // CHECK: nvvm.griddepcontrol launch_dependents511 nvvm.griddepcontrol launch_dependents512 return513}514 515// CHECK-LABEL: @mapa516func.func @mapa(%a: !llvm.ptr, %a_shared: !llvm.ptr<3>, %b : i32) {517 // CHECK: nvvm.mapa %{{.*}}518 %0 = nvvm.mapa %a, %b: !llvm.ptr -> !llvm.ptr519 // CHECK: nvvm.mapa %{{.*}}520 %1 = nvvm.mapa %a_shared, %b: !llvm.ptr<3> -> !llvm.ptr<7>521 return522}523 524// CHECK-LABEL: @match_sync525func.func @match_sync(%val32: i32, %val64: i64, %thread_mask: i32) {526 // CHECK: nvvm.match.sync any %{{.*}}, %{{.*}} : i32 -> i32527 %0 = nvvm.match.sync any %thread_mask, %val32 : i32 -> i32528 // CHECK: nvvm.match.sync all %{{.*}}, %{{.*}} : i32 -> !llvm.struct<(i32, i1)>529 %1 = nvvm.match.sync all %thread_mask, %val32 : i32 -> !llvm.struct<(i32, i1)>530 // CHECK: nvvm.match.sync any %{{.*}}, %{{.*}} : i64 -> i32531 %2 = nvvm.match.sync any %thread_mask, %val64 : i64 -> i32532 // CHECK: nvvm.match.sync all %{{.*}}, %{{.*}} : i64 -> !llvm.struct<(i32, i1)>533 %3 = nvvm.match.sync all %thread_mask, %val64 : i64 -> !llvm.struct<(i32, i1)>534 return 535}536 537// CHECK-LABEL: @st_bulk538func.func @st_bulk(%addr_gen: !llvm.ptr, %addr_shared: !llvm.ptr<3>, %size: i64) {539 // CHECK: nvvm.st.bulk %{{.*}}, size = %{{.*}} : !llvm.ptr540 nvvm.st.bulk %addr_gen, size = %size, init = 0 : !llvm.ptr541 // CHECK: nvvm.st.bulk %{{.*}}, size = %{{.*}} : !llvm.ptr<3>542 nvvm.st.bulk %addr_shared, size = %size, init = 0 : !llvm.ptr<3>543 return544}545 546// CHECK-LABEL: @dot_accumulate_4way547func.func @dot_accumulate_4way(%a_vec: vector<4xi8>, %b_vec: vector<4xi8>, %c: i32) {548 // CHECK: nvvm.dot.accumulate.4way %{{.*}}, %{{.*}}, %{{.*}} : vector<4xi8>, vector<4xi8>549 %1 = nvvm.dot.accumulate.4way %a_vec <unsigned>, %b_vec <unsigned>, %c: vector<4xi8>, vector<4xi8>550 // CHECK: nvvm.dot.accumulate.4way %{{.*}}, %{{.*}}, %{{.*}} : vector<4xi8>, vector<4xi8>551 %3 = nvvm.dot.accumulate.4way %a_vec <signed>, %b_vec <signed>, %c: vector<4xi8>, vector<4xi8>552 return553}554 555// CHECK-LABEL: @dot_accumulate_2way556func.func @dot_accumulate_2way(%a_vec: vector<2xi16>, %b_vec: vector<4xi8>, %c: i32) {557 // CHECK: nvvm.dot.accumulate.2way %{{.*}}, %{{.*}}, %{{.*}} {b_hi = false} : vector<2xi16>, vector<4xi8>558 %1 = nvvm.dot.accumulate.2way %a_vec <unsigned>, %b_vec <unsigned>, %c {b_hi = false}: vector<2xi16>, vector<4xi8>559 // CHECK: nvvm.dot.accumulate.2way %{{.*}}, %{{.*}}, %{{.*}} {b_hi = true} : vector<2xi16>, vector<4xi8>560 %3 = nvvm.dot.accumulate.2way %a_vec <signed>, %b_vec <signed>, %c {b_hi = true}: vector<2xi16>, vector<4xi8>561 return562}563 564// CHECK-LABEL: @prefetch565func.func @prefetch(%gen_ptr: !llvm.ptr, %local_ptr: !llvm.ptr<5>, %global_ptr: !llvm.ptr<1>, %const_ptr: !llvm.ptr<4>) {566 // CHECK: nvvm.prefetch level = L1, %{{.*}}567 nvvm.prefetch level = L1, %gen_ptr : !llvm.ptr<0>568 // CHECK: nvvm.prefetch level = L1, %{{.*}}569 nvvm.prefetch level = L1, %local_ptr : !llvm.ptr<5>570 // CHECK: nvvm.prefetch level = L1, %{{.*}}571 nvvm.prefetch level = L1, %global_ptr : !llvm.ptr<1>572 // CHECK: nvvm.prefetch level = L2, %{{.*}}573 nvvm.prefetch level = L2, %gen_ptr : !llvm.ptr<0>574 // CHECK: nvvm.prefetch level = L2, %{{.*}}575 nvvm.prefetch level = L2, %local_ptr : !llvm.ptr<5>576 // CHECK: nvvm.prefetch level = L2, %{{.*}}577 nvvm.prefetch level = L2, %global_ptr : !llvm.ptr<1>578 // CHECK: nvvm.prefetch level = L2, evict_priority = evict_last, %{{.*}}579 nvvm.prefetch level = L2, evict_priority = evict_last, %global_ptr :580 !llvm.ptr<1>581 // CHECK: nvvm.prefetch level = L2, evict_priority = evict_normal, %{{.*}}582 nvvm.prefetch level = L2, evict_priority = evict_normal, %global_ptr : !llvm.ptr<1>583 // CHECK: nvvm.prefetch level = L1 uniform, %{{.*}}584 nvvm.prefetch level = L1 uniform, %gen_ptr : !llvm.ptr585 // CHECK: nvvm.prefetch tensormap, %{{.*}}586 nvvm.prefetch tensormap, %gen_ptr : !llvm.ptr587 // CHECK: nvvm.prefetch tensormap, %{{.*}}588 nvvm.prefetch tensormap, %const_ptr : !llvm.ptr<4>589 // CHECK: nvvm.prefetch tensormap in_param_space, %{{.*}}590 nvvm.prefetch tensormap in_param_space, %gen_ptr : !llvm.ptr591 return592}593 594// CHECK-LABEL: @prefetch_tensormap595func.func @prefetch_tensormap(%gen_ptr: !llvm.ptr, %const_ptr: !llvm.ptr<4>) {596 return597}598 599// CHECK-LABEL: @nvvm_address_space600func.func private @nvvm_address_space(601 !ptr.ptr<#nvvm.memory_space<global>>,602 !ptr.ptr<#nvvm.memory_space<shared>>,603 !ptr.ptr<#nvvm.memory_space<constant>>,604 !ptr.ptr<#nvvm.memory_space<local>>,605 !ptr.ptr<#nvvm.memory_space<tensor>>,606 !ptr.ptr<#nvvm.memory_space<shared_cluster>>607 ) -> !ptr.ptr<#nvvm.memory_space<generic>>608 609// -----610 611// Just check these don't emit errors.612gpu.module @module_1 [#nvvm.target<chip = "sm_90", features = "+ptx70", link = ["my_device_lib.bc"], flags = {fast, ftz}>] {613}614 615gpu.module @module_2 [#nvvm.target<chip = "sm_90">, #nvvm.target<chip = "sm_80">, #nvvm.target<chip = "sm_70">] {616}617 618// CHECK-LABEL: nvvm.grid_constant619llvm.func @kernel_func(%arg0: !llvm.ptr {llvm.byval = i32, nvvm.grid_constant}) attributes {nvvm.kernel} {620 llvm.return621}622 623// -----624 625// expected-error @below {{'"nvvm.grid_constant"' attribute must be present only on kernel arguments}}626llvm.func @kernel_func(%arg0: !llvm.ptr {llvm.byval = i32, nvvm.grid_constant}) {627 llvm.return628}629 630// -----631 632// expected-error @below {{'"nvvm.grid_constant"' attribute requires the argument to also have attribute 'llvm.byval'}}633llvm.func @kernel_func(%arg0: !llvm.ptr {nvvm.grid_constant}) attributes {nvvm.kernel} {634 llvm.return635}636 637// -----638 639// expected-error @below {{'"nvvm.grid_constant"' must be a unit attribute}}640llvm.func @kernel_func(%arg0: !llvm.ptr {llvm.byval = i32, nvvm.grid_constant = true}) attributes {nvvm.kernel} {641 llvm.return642}643