1// RUN: mlir-opt %s -convert-gpu-to-nvvm='has-redux=1' -split-input-file | FileCheck %s 2// RUN: mlir-opt %s -convert-gpu-to-nvvm='has-redux=1 use-bare-ptr-memref-call-conv=1' -split-input-file | FileCheck %s --check-prefix=CHECK-BARE 3// RUN: mlir-opt %s -transform-interpreter | FileCheck %s 4 5gpu.module @test_module_0 { 6 // CHECK-LABEL: func @gpu_index_ops() 7 func.func @gpu_index_ops() 8 -> (index, index, index, index, index, index, 9 index, index, index, index, index, index, 10 index) { 11 12 // CHECK: = nvvm.read.ptx.sreg.tid.x : i32 13 // CHECK: = llvm.sext %{{.*}} : i32 to i64 14 %tIdX = gpu.thread_id x 15 // CHECK: = nvvm.read.ptx.sreg.tid.y : i32 16 // CHECK: = llvm.sext %{{.*}} : i32 to i64 17 %tIdY = gpu.thread_id y 18 // CHECK: = nvvm.read.ptx.sreg.tid.z : i32 19 // CHECK: = llvm.sext %{{.*}} : i32 to i64 20 %tIdZ = gpu.thread_id z 21 22 // CHECK: = nvvm.read.ptx.sreg.ntid.x : i32 23 // CHECK: = llvm.sext %{{.*}} : i32 to i64 24 %bDimX = gpu.block_dim x 25 // CHECK: = nvvm.read.ptx.sreg.ntid.y : i32 26 // CHECK: = llvm.sext %{{.*}} : i32 to i64 27 %bDimY = gpu.block_dim y 28 // CHECK: = nvvm.read.ptx.sreg.ntid.z : i32 29 // CHECK: = llvm.sext %{{.*}} : i32 to i64 30 %bDimZ = gpu.block_dim z 31 32 // CHECK: = nvvm.read.ptx.sreg.ctaid.x : i32 33 // CHECK: = llvm.sext %{{.*}} : i32 to i64 34 %bIdX = gpu.block_id x 35 // CHECK: = nvvm.read.ptx.sreg.ctaid.y : i32 36 // CHECK: = llvm.sext %{{.*}} : i32 to i64 37 %bIdY = gpu.block_id y 38 // CHECK: = nvvm.read.ptx.sreg.ctaid.z : i32 39 // CHECK: = llvm.sext %{{.*}} : i32 to i64 40 %bIdZ = gpu.block_id z 41 42 // CHECK: = nvvm.read.ptx.sreg.nctaid.x : i32 43 // CHECK: = llvm.sext %{{.*}} : i32 to i64 44 %gDimX = gpu.grid_dim x 45 // CHECK: = nvvm.read.ptx.sreg.nctaid.y : i32 46 // CHECK: = llvm.sext %{{.*}} : i32 to i64 47 %gDimY = gpu.grid_dim y 48 // CHECK: = nvvm.read.ptx.sreg.nctaid.z : i32 49 // CHECK: = llvm.sext %{{.*}} : i32 to i64 50 %gDimZ = gpu.grid_dim z 51 52 53 // CHECK: = nvvm.read.ptx.sreg.laneid range <i32, 0, 32> : i32 54 // CHECK: = llvm.sext %{{.*}} : i32 to i64 55 %laneId = gpu.lane_id 56 57 func.return %tIdX, %tIdY, %tIdZ, %bDimX, %bDimY, %bDimZ, 58 %bIdX, %bIdY, %bIdZ, %gDimX, %gDimY, %gDimZ, 59 %laneId 60 : index, index, index, index, index, index, 61 index, index, index, index, index, index, 62 index 63 } 64} 65 66 67 68gpu.module @test_module_1 { 69 // CHECK-LABEL: func @gpu_index_comp 70 func.func @gpu_index_comp(%idx : index) -> index { 71 // CHECK: = llvm.add %{{.*}}, %{{.*}} : i64 72 %0 = arith.addi %idx, %idx : index 73 // CHECK: llvm.return %{{.*}} : i64 74 func.return %0 : index 75 } 76} 77 78 79 80gpu.module @test_module_2 { 81 // CHECK-LABEL: func @gpu_all_reduce_op() 82 gpu.func @gpu_all_reduce_op() { 83 %arg0 = arith.constant 1.0 : f32 84 // TODO: Check full IR expansion once lowering has settled. 85 // CHECK: nvvm.shfl.sync bfly {{.*}} 86 // CHECK: nvvm.barrier0 87 // CHECK: llvm.fadd 88 %result = gpu.all_reduce add %arg0 uniform {} : (f32) -> (f32) 89 90 gpu.return 91 } 92} 93 94 95 96gpu.module @test_module_3 { 97 // CHECK-LABEL: func @gpu_all_reduce_region() 98 gpu.func @gpu_all_reduce_region() { 99 %arg0 = arith.constant 1 : i32 100 // TODO: Check full IR expansion once lowering has settled. 101 // CHECK: nvvm.shfl.sync bfly {{.*}} 102 // CHECK: nvvm.barrier0 103 %result = gpu.all_reduce %arg0 uniform { 104 ^bb(%lhs : i32, %rhs : i32): 105 %xor = arith.xori %lhs, %rhs : i32 106 "gpu.yield"(%xor) : (i32) -> () 107 } : (i32) -> (i32) 108 gpu.return 109 } 110} 111 112 113 114gpu.module @test_module_4 { 115 // CHECK-LABEL: func @gpu_shuffle() 116 func.func @gpu_shuffle() -> (f32, f32, f32, f32, i1, i1, i1, i1) { 117 // CHECK: %[[#VALUE:]] = llvm.mlir.constant(1.000000e+00 : f32) : f32 118 %arg0 = arith.constant 1.0 : f32 119 // CHECK: %[[#OFFSET:]] = llvm.mlir.constant(4 : i32) : i32 120 %arg1 = arith.constant 4 : i32 121 // CHECK: %[[#WIDTH:]] = llvm.mlir.constant(23 : i32) : i32 122 %arg2 = arith.constant 23 : i32 123 // CHECK: %[[#ONE:]] = llvm.mlir.constant(1 : i32) : i32 124 // CHECK: %[[#MINUS_ONE:]] = llvm.mlir.constant(-1 : i32) : i32 125 // CHECK: %[[#THIRTY_TWO:]] = llvm.mlir.constant(32 : i32) : i32 126 // CHECK: %[[#NUM_LANES:]] = llvm.sub %[[#THIRTY_TWO]], %[[#WIDTH]] : i32 127 // CHECK: %[[#MASK:]] = llvm.lshr %[[#MINUS_ONE]], %[[#NUM_LANES]] : i32 128 // CHECK: %[[#CLAMP:]] = llvm.sub %[[#WIDTH]], %[[#ONE]] : i32 129 // CHECK: %[[#SHFL:]] = nvvm.shfl.sync bfly %[[#MASK]], %[[#VALUE]], %[[#OFFSET]], %[[#CLAMP]] {return_value_and_is_valid} : f32 -> !llvm.struct<(f32, i1)> 130 // CHECK: llvm.extractvalue %[[#SHFL]][0] : !llvm.struct<(f32, i1)> 131 // CHECK: llvm.extractvalue %[[#SHFL]][1] : !llvm.struct<(f32, i1)> 132 %shfl, %pred = gpu.shuffle xor %arg0, %arg1, %arg2 : f32 133 // CHECK: %[[#ONE:]] = llvm.mlir.constant(1 : i32) : i32 134 // CHECK: %[[#MINUS_ONE:]] = llvm.mlir.constant(-1 : i32) : i32 135 // CHECK: %[[#THIRTY_TWO:]] = llvm.mlir.constant(32 : i32) : i32 136 // CHECK: %[[#NUM_LANES:]] = llvm.sub %[[#THIRTY_TWO]], %[[#WIDTH]] : i32 137 // CHECK: %[[#MASK:]] = llvm.lshr %[[#MINUS_ONE]], %[[#NUM_LANES]] : i32 138 // CHECK: %[[#SHFL:]] = nvvm.shfl.sync up %[[#MASK]], %[[#VALUE]], %[[#OFFSET]], %[[#NUM_LANES]] {return_value_and_is_valid} : f32 -> !llvm.struct<(f32, i1)> 139 // CHECK: llvm.extractvalue %[[#SHFL]][0] : !llvm.struct<(f32, i1)> 140 // CHECK: llvm.extractvalue %[[#SHFL]][1] : !llvm.struct<(f32, i1)> 141 %shflu, %predu = gpu.shuffle up %arg0, %arg1, %arg2 : f32 142 // CHECK: nvvm.shfl.sync down {{.*}} {return_value_and_is_valid} : f32 -> !llvm.struct<(f32, i1)> 143 %shfld, %predd = gpu.shuffle down %arg0, %arg1, %arg2 : f32 144 // CHECK: nvvm.shfl.sync idx {{.*}} {return_value_and_is_valid} : f32 -> !llvm.struct<(f32, i1)> 145 %shfli, %predi = gpu.shuffle idx %arg0, %arg1, %arg2 : f32 146 147 func.return %shfl, %shflu, %shfld, %shfli, %pred, %predu, %predd, %predi 148 : f32, f32,f32, f32, i1, i1, i1, i1 149 } 150 151 // CHECK-LABEL: func @gpu_shuffle_unused_pred() 152 func.func @gpu_shuffle_unused_pred() -> (f32, f32, f32, f32) { 153 // CHECK: %[[#VALUE:]] = llvm.mlir.constant(1.000000e+00 : f32) : f32 154 %arg0 = arith.constant 1.0 : f32 155 // CHECK: %[[#OFFSET:]] = llvm.mlir.constant(4 : i32) : i32 156 %arg1 = arith.constant 4 : i32 157 // CHECK: %[[#WIDTH:]] = llvm.mlir.constant(23 : i32) : i32 158 %arg2 = arith.constant 23 : i32 159 // CHECK: %[[#ONE:]] = llvm.mlir.constant(1 : i32) : i32 160 // CHECK: %[[#MINUS_ONE:]] = llvm.mlir.constant(-1 : i32) : i32 161 // CHECK: %[[#THIRTY_TWO:]] = llvm.mlir.constant(32 : i32) : i32 162 // CHECK: %[[#NUM_LANES:]] = llvm.sub %[[#THIRTY_TWO]], %[[#WIDTH]] : i32 163 // CHECK: %[[#MASK:]] = llvm.lshr %[[#MINUS_ONE]], %[[#NUM_LANES]] : i32 164 // CHECK: %[[#CLAMP:]] = llvm.sub %[[#WIDTH]], %[[#ONE]] : i32 165 // CHECK: %[[#SHFL:]] = nvvm.shfl.sync bfly %[[#MASK]], %[[#VALUE]], %[[#OFFSET]], %[[#CLAMP]] : f32 -> f32 166 %shfl, %pred = gpu.shuffle xor %arg0, %arg1, %arg2 : f32 167 // CHECK: %[[#ONE:]] = llvm.mlir.constant(1 : i32) : i32 168 // CHECK: %[[#MINUS_ONE:]] = llvm.mlir.constant(-1 : i32) : i32 169 // CHECK: %[[#THIRTY_TWO:]] = llvm.mlir.constant(32 : i32) : i32 170 // CHECK: %[[#NUM_LANES:]] = llvm.sub %[[#THIRTY_TWO]], %[[#WIDTH]] : i32 171 // CHECK: %[[#MASK:]] = llvm.lshr %[[#MINUS_ONE]], %[[#NUM_LANES]] : i32 172 // CHECK: %[[#SHFL:]] = nvvm.shfl.sync up %[[#MASK]], %[[#VALUE]], %[[#OFFSET]], %[[#NUM_LANES]] : f32 -> f32 173 %shflu, %predu = gpu.shuffle up %arg0, %arg1, %arg2 : f32 174 // CHECK: nvvm.shfl.sync down {{.*}} : f32 -> f32 175 %shfld, %predd = gpu.shuffle down %arg0, %arg1, %arg2 : f32 176 // CHECK: nvvm.shfl.sync idx {{.*}} : f32 -> f32 177 %shfli, %predi = gpu.shuffle idx %arg0, %arg1, %arg2 : f32 178 179 func.return %shfl, %shflu, %shfld, %shfli : f32, f32,f32, f32 180 } 181} 182 183gpu.module @test_module_5 { 184 // CHECK-LABEL: func @gpu_sync() 185 func.func @gpu_sync() { 186 // CHECK: nvvm.barrier0 187 gpu.barrier 188 func.return 189 } 190} 191 192 193 194gpu.module @test_module_6 { 195 // CHECK: llvm.func @__nv_fabsf(f32) -> f32 196 // CHECK: llvm.func @__nv_fabs(f64) -> f64 197 // CHECK-LABEL: func @gpu_fabs 198 func.func @gpu_fabs(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) { 199 %result32 = math.absf %arg_f32 : f32 200 // CHECK: llvm.call @__nv_fabsf(%{{.*}}) : (f32) -> f32 201 %result64 = math.absf %arg_f64 : f64 202 // CHECK: llvm.call @__nv_fabs(%{{.*}}) : (f64) -> f64 203 func.return %result32, %result64 : f32, f64 204 } 205} 206 207 208 209gpu.module @test_module_7 { 210 // CHECK: llvm.func @__nv_cbrtf(f32) -> f32 211 // CHECK: llvm.func @__nv_cbrt(f64) -> f64 212 // CHECK-LABEL: func @gpu_cbrt 213 func.func @gpu_cbrt(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) { 214 %result32 = math.cbrt %arg_f32 : f32 215 // CHECK: llvm.call @__nv_cbrtf(%{{.*}}) : (f32) -> f32 216 %result64 = math.cbrt %arg_f64 : f64 217 // CHECK: llvm.call @__nv_cbrt(%{{.*}}) : (f64) -> f64 218 func.return %result32, %result64 : f32, f64 219 } 220} 221 222 223 224gpu.module @test_module_8 { 225 // CHECK: llvm.func @__nv_ceilf(f32) -> f32 226 // CHECK: llvm.func @__nv_ceil(f64) -> f64 227 // CHECK-LABEL: func @gpu_ceil 228 func.func @gpu_ceil(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) { 229 %result32 = math.ceil %arg_f32 : f32 230 // CHECK: llvm.call @__nv_ceilf(%{{.*}}) : (f32) -> f32 231 %result64 = math.ceil %arg_f64 : f64 232 // CHECK: llvm.call @__nv_ceil(%{{.*}}) : (f64) -> f64 233 func.return %result32, %result64 : f32, f64 234 } 235} 236 237 238 239gpu.module @test_module_9 { 240 // CHECK: llvm.func @__nv_floorf(f32) -> f32 241 // CHECK: llvm.func @__nv_floor(f64) -> f64 242 // CHECK-LABEL: func @gpu_floor 243 func.func @gpu_floor(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) { 244 %result32 = math.floor %arg_f32 : f32 245 // CHECK: llvm.call @__nv_floorf(%{{.*}}) : (f32) -> f32 246 %result64 = math.floor %arg_f64 : f64 247 // CHECK: llvm.call @__nv_floor(%{{.*}}) : (f64) -> f64 248 func.return %result32, %result64 : f32, f64 249 } 250} 251 252 253 254gpu.module @test_module_10 { 255 // CHECK: llvm.func @__nv_cosf(f32) -> f32 256 // CHECK: llvm.func @__nv_cos(f64) -> f64 257 // CHECK: llvm.func @__nv_fast_cosf(f32) -> f32 258 // CHECK-LABEL: func @gpu_cos 259 func.func @gpu_cos(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64, f32) { 260 %result32 = math.cos %arg_f32 : f32 261 // CHECK: llvm.call @__nv_cosf(%{{.*}}) : (f32) -> f32 262 %result64 = math.cos %arg_f64 : f64 263 // CHECK: llvm.call @__nv_cos(%{{.*}}) : (f64) -> f64 264 %result32Fast = math.cos %arg_f32 fastmath<afn> : f32 265 // CHECK: llvm.call @__nv_fast_cosf(%{{.*}}) : (f32) -> f32 266 func.return %result32, %result64, %result32Fast : f32, f64, f32 267 } 268} 269 270 271gpu.module @test_module_11 { 272 // CHECK: llvm.func @__nv_expf(f32) -> f32 273 // CHECK: llvm.func @__nv_exp(f64) -> f64 274 // CHECK: llvm.func @__nv_fast_expf(f32) -> f32 275 // CHECK-LABEL: func @gpu_exp 276 func.func @gpu_exp(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64, f32) { 277 %result32 = math.exp %arg_f32 : f32 278 // CHECK: llvm.call @__nv_expf(%{{.*}}) : (f32) -> f32 279 %result64 = math.exp %arg_f64 : f64 280 // CHECK: llvm.call @__nv_exp(%{{.*}}) : (f64) -> f64 281 %result32Fast = math.exp %arg_f32 fastmath<fast> : f32 282 // CHECK: llvm.call @__nv_fast_expf(%{{.*}}) : (f32) -> f32 283 func.return %result32, %result64, %result32Fast : f32, f64, f32 284 } 285} 286 287 288gpu.module @test_module_12 { 289 // CHECK: llvm.func @__nv_exp2f(f32) -> f32 290 // CHECK: llvm.func @__nv_exp2(f64) -> f64 291 // CHECK-LABEL: func @gpu_exp2 292 func.func @gpu_exp2(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) { 293 %result32 = math.exp2 %arg_f32 : f32 294 // CHECK: llvm.call @__nv_exp2f(%{{.*}}) : (f32) -> f32 295 %result64 = math.exp2 %arg_f64 : f64 296 // CHECK: llvm.call @__nv_exp2(%{{.*}}) : (f64) -> f64 297 func.return %result32, %result64 : f32, f64 298 } 299} 300 301 302 303gpu.module @test_module_13 { 304 // CHECK: llvm.func @__nv_logf(f32) -> f32 305 // CHECK: llvm.func @__nv_log(f64) -> f64 306 // CHECK: llvm.func @__nv_fast_logf(f32) -> f32 307 // CHECK-LABEL: func @gpu_log 308 func.func @gpu_log(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64, f32) { 309 %result32 = math.log %arg_f32 : f32 310 // CHECK: llvm.call @__nv_logf(%{{.*}}) : (f32) -> f32 311 %result64 = math.log %arg_f64 : f64 312 // CHECK: llvm.call @__nv_log(%{{.*}}) : (f64) -> f64 313 %result32Fast = math.log %arg_f32 fastmath<afn> : f32 314 // CHECK: llvm.call @__nv_fast_logf(%{{.*}}) : (f32) -> f32 315 func.return %result32, %result64, %result32Fast : f32, f64, f32 316 } 317} 318 319 320 321gpu.module @test_module_14 { 322 // CHECK: llvm.func @__nv_log10f(f32) -> f32 323 // CHECK: llvm.func @__nv_log10(f64) -> f64 324 // CHECK: llvm.func @__nv_fast_log10f(f32) -> f32 325 // CHECK-LABEL: func @gpu_log10 326 func.func @gpu_log10(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64, f32) { 327 %result32 = math.log10 %arg_f32 : f32 328 // CHECK: llvm.call @__nv_log10f(%{{.*}}) : (f32) -> f32 329 %result64 = math.log10 %arg_f64 : f64 330 // CHECK: llvm.call @__nv_log10(%{{.*}}) : (f64) -> f64 331 %result32Fast = math.log10 %arg_f32 fastmath<afn> : f32 332 // CHECK: llvm.call @__nv_fast_log10f(%{{.*}}) : (f32) -> f32 333 func.return %result32, %result64, %result32Fast : f32, f64, f32 334 } 335} 336 337 338 339gpu.module @test_module_15 { 340 // CHECK: llvm.func @__nv_log1pf(f32) -> f32 341 // CHECK: llvm.func @__nv_log1p(f64) -> f64 342 // CHECK-LABEL: func @gpu_log1p 343 func.func @gpu_log1p(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) { 344 %result32 = math.log1p %arg_f32 : f32 345 // CHECK: llvm.call @__nv_log1pf(%{{.*}}) : (f32) -> f32 346 %result64 = math.log1p %arg_f64 : f64 347 // CHECK: llvm.call @__nv_log1p(%{{.*}}) : (f64) -> f64 348 func.return %result32, %result64 : f32, f64 349 } 350} 351 352 353 354gpu.module @test_module_16 { 355 // CHECK: llvm.func @__nv_log2f(f32) -> f32 356 // CHECK: llvm.func @__nv_log2(f64) -> f64 357 // CHECK: llvm.func @__nv_fast_log2f(f32) -> f32 358 // CHECK-LABEL: func @gpu_log2 359 func.func @gpu_log2(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64, f32) { 360 %result32 = math.log2 %arg_f32 : f32 361 // CHECK: llvm.call @__nv_log2f(%{{.*}}) : (f32) -> f32 362 %result64 = math.log2 %arg_f64 : f64 363 // CHECK: llvm.call @__nv_log2(%{{.*}}) : (f64) -> f64 364 %result32Fast = math.log2 %arg_f32 fastmath<fast> : f32 365 // CHECK: llvm.call @__nv_fast_log2f(%{{.*}}) : (f32) -> f32 366 func.return %result32, %result64, %result32Fast : f32, f64, f32 367 } 368} 369 370 371 372gpu.module @test_module_17 { 373 // CHECK: llvm.func @__nv_sinf(f32) -> f32 374 // CHECK: llvm.func @__nv_sin(f64) -> f64 375 // CHECK: llvm.func @__nv_fast_sinf(f32) -> f32 376 // CHECK-LABEL: func @gpu_sin 377 func.func @gpu_sin(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64, f32) { 378 %result32 = math.sin %arg_f32 : f32 379 // CHECK: llvm.call @__nv_sinf(%{{.*}}) : (f32) -> f32 380 %result64 = math.sin %arg_f64 : f64 381 // CHECK: llvm.call @__nv_sin(%{{.*}}) : (f64) -> f64 382 %result32Fast = math.sin %arg_f32 fastmath<fast> : f32 383 // CHECK: llvm.call @__nv_fast_sinf(%{{.*}}) : (f32) -> f32 384 func.return %result32, %result64, %result32Fast : f32, f64, f32 385 } 386} 387 388 389 390gpu.module @test_module_18 { 391 // CHECK: llvm.func @__nv_tanf(f32) -> f32 392 // CHECK: llvm.func @__nv_tan(f64) -> f64 393 // CHECK: llvm.func @__nv_fast_tanf(f32) -> f32 394 // CHECK-LABEL: func @gpu_tan 395 func.func @gpu_tan(%arg_f16 : f16, %arg_f32 : f32, %arg_f64 : f64) -> (f16, f32, f64, f32) { 396 %result16 = math.tan %arg_f16 : f16 397 // CHECK: llvm.fpext %{{.*}} : f16 to f32 398 // CHECK-NEXT: llvm.call @__nv_tanf(%{{.*}}) : (f32) -> f32 399 // CHECK-NEXT: llvm.fptrunc %{{.*}} : f32 to f16 400 %result32 = math.tan %arg_f32 : f32 401 // CHECK: llvm.call @__nv_tanf(%{{.*}}) : (f32) -> f32 402 %result64 = math.tan %arg_f64 : f64 403 // CHECK: llvm.call @__nv_tan(%{{.*}}) : (f64) -> f64 404 %result32Fast = math.tan %arg_f32 fastmath<fast> : f32 405 // CHECK: llvm.call @__nv_fast_tanf(%{{.*}}) : (f32) -> f32 406 func.return %result16, %result32, %result64, %result32Fast : f16, f32, f64, f32 407 } 408} 409 410 411 412gpu.module @test_module_19 { 413 // CHECK: llvm.func @__nv_tanhf(f32) -> f32 414 // CHECK: llvm.func @__nv_tanh(f64) -> f64 415 // CHECK-LABEL: func @gpu_tanh 416 func.func @gpu_tanh(%arg_f16 : f16, %arg_f32 : f32, %arg_f64 : f64) -> (f16, f32, f64) { 417 %result16 = math.tanh %arg_f16 : f16 418 // CHECK: llvm.fpext %{{.*}} : f16 to f32 419 // CHECK-NEXT: llvm.call @__nv_tanhf(%{{.*}}) : (f32) -> f32 420 // CHECK-NEXT: llvm.fptrunc %{{.*}} : f32 to f16 421 %result32 = math.tanh %arg_f32 : f32 422 // CHECK: llvm.call @__nv_tanhf(%{{.*}}) : (f32) -> f32 423 %result64 = math.tanh %arg_f64 : f64 424 // CHECK: llvm.call @__nv_tanh(%{{.*}}) : (f64) -> f64 425 func.return %result16, %result32, %result64 : f16, f32, f64 426 } 427} 428 429 430 431gpu.module @test_module_20 { 432 // CHECK: llvm.func @__nv_rsqrtf(f32) -> f32 433 // CHECK: llvm.func @__nv_rsqrt(f64) -> f64 434 // CHECK-LABEL: func @gpu_rsqrt 435 func.func @gpu_rsqrt(%arg_f16 : f16, %arg_f32 : f32, %arg_f64 : f64) 436 -> (f16, f32, f64) { 437 %result16 = math.rsqrt %arg_f16 : f16 438 // CHECK: llvm.fpext %{{.*}} : f16 to f32 439 // CHECK-NEXT: llvm.call @__nv_rsqrtf(%{{.*}}) : (f32) -> f32 440 // CHECK-NEXT: llvm.fptrunc %{{.*}} : f32 to f16 441 %result32 = math.rsqrt %arg_f32 : f32 442 // CHECK: llvm.call @__nv_rsqrtf(%{{.*}}) : (f32) -> f32 443 %result64 = math.rsqrt %arg_f64 : f64 444 // CHECK: llvm.call @__nv_rsqrt(%{{.*}}) : (f64) -> f64 445 func.return %result16, %result32, %result64 : f16, f32, f64 446 } 447} 448 449 450 451gpu.module @test_module_21 { 452 // CHECK: llvm.func @__nv_sqrtf(f32) -> f32 453 // CHECK: llvm.func @__nv_sqrt(f64) -> f64 454 // CHECK-LABEL: func @gpu_sqrt 455 func.func @gpu_sqrt(%arg_f16 : f16, %arg_f32 : f32, %arg_f64 : f64) 456 -> (f16, f32, f64) { 457 %result16 = math.sqrt %arg_f16 : f16 458 // CHECK: llvm.fpext %{{.*}} : f16 to f32 459 // CHECK-NEXT: llvm.call @__nv_sqrtf(%{{.*}}) : (f32) -> f32 460 // CHECK-NEXT: llvm.fptrunc %{{.*}} : f32 to f16 461 %result32 = math.sqrt %arg_f32 : f32 462 // CHECK: llvm.call @__nv_sqrtf(%{{.*}}) : (f32) -> f32 463 %result64 = math.sqrt %arg_f64 : f64 464 // CHECK: llvm.call @__nv_sqrt(%{{.*}}) : (f64) -> f64 465 func.return %result16, %result32, %result64 : f16, f32, f64 466 } 467} 468 469 470 471gpu.module @test_module_22 { 472 // CHECK: llvm.func @__nv_atanf(f32) -> f32 473 // CHECK: llvm.func @__nv_atan(f64) -> f64 474 // CHECK-LABEL: func @gpu_atan 475 func.func @gpu_atan(%arg_f16 : f16, %arg_f32 : f32, %arg_f64 : f64) 476 -> (f16, f32, f64) { 477 %result16 = math.atan %arg_f16 : f16 478 // CHECK: llvm.fpext %{{.*}} : f16 to f32 479 // CHECK-NEXT: llvm.call @__nv_atanf(%{{.*}}) : (f32) -> f32 480 // CHECK-NEXT: llvm.fptrunc %{{.*}} : f32 to f16 481 %result32 = math.atan %arg_f32 : f32 482 // CHECK: llvm.call @__nv_atanf(%{{.*}}) : (f32) -> f32 483 %result64 = math.atan %arg_f64 : f64 484 // CHECK: llvm.call @__nv_atan(%{{.*}}) : (f64) -> f64 485 func.return %result16, %result32, %result64 : f16, f32, f64 486 } 487} 488 489 490 491gpu.module @test_module_23 { 492 // CHECK: llvm.func @__nv_atan2f(f32, f32) -> f32 493 // CHECK: llvm.func @__nv_atan2(f64, f64) -> f64 494 // CHECK-LABEL: func @gpu_atan2 495 func.func @gpu_atan2(%arg_f16 : f16, %arg_f32 : f32, %arg_f64 : f64) 496 -> (f16, f32, f64) { 497 %result16 = math.atan2 %arg_f16, %arg_f16 : f16 498 // CHECK: llvm.fpext %{{.*}} : f16 to f32 499 // CHECK: llvm.fpext %{{.*}} : f16 to f32 500 // CHECK-NEXT: llvm.call @__nv_atan2f(%{{.*}}) : (f32, f32) -> f32 501 // CHECK-NEXT: llvm.fptrunc %{{.*}} : f32 to f16 502 %result32 = math.atan2 %arg_f32, %arg_f32 : f32 503 // CHECK: llvm.call @__nv_atan2f(%{{.*}}) : (f32, f32) -> f32 504 %result64 = math.atan2 %arg_f64, %arg_f64 : f64 505 // CHECK: llvm.call @__nv_atan2(%{{.*}}) : (f64, f64) -> f64 506 func.return %result16, %result32, %result64 : f16, f32, f64 507 } 508} 509 510 511 512// Test that we handled properly operation with SymbolTable other than module op 513gpu.module @test_module_24 { 514 "test.symbol_scope"() ({ 515 // CHECK: test.symbol_scope 516 // CHECK: llvm.func @__nv_expf(f32) -> f32 517 // CHECK: llvm.func @__nv_exp(f64) -> f64 518 // CHECK: llvm.func @__nv_fast_expf(f32) -> f32 519 // CHECK-LABEL: func @gpu_exp 520 func.func @gpu_exp(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64, f32) { 521 %result32 = math.exp %arg_f32 : f32 522 // CHECK: llvm.call @__nv_expf(%{{.*}}) : (f32) -> f32 523 %result64 = math.exp %arg_f64 : f64 524 // CHECK: llvm.call @__nv_exp(%{{.*}}) : (f64) -> f64 525 %result32Fast = math.exp %arg_f32 fastmath<afn> : f32 526 // CHECK: llvm.call @__nv_fast_expf(%{{.*}}) : (f32) -> f32 527 func.return %result32, %result64, %result32Fast : f32, f64, f32 528 } 529 "test.finish" () : () -> () 530 }) : () -> () 531} 532 533 534 535gpu.module @test_module_25 { 536 // CHECK: llvm.func @__nv_expm1f(f32) -> f32 537 // CHECK: llvm.func @__nv_expm1(f64) -> f64 538 // CHECK-LABEL: func @gpu_expm1 539 func.func @gpu_expm1(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) { 540 %result32 = math.expm1 %arg_f32 : f32 541 // CHECK: llvm.call @__nv_expm1f(%{{.*}}) : (f32) -> f32 542 %result64 = math.expm1 %arg_f64 : f64 543 // CHECK: llvm.call @__nv_expm1(%{{.*}}) : (f64) -> f64 544 func.return %result32, %result64 : f32, f64 545 } 546} 547 548 549 550gpu.module @test_module_26 { 551 // CHECK: llvm.func @__nv_powf(f32, f32) -> f32 552 // CHECK: llvm.func @__nv_pow(f64, f64) -> f64 553 // CHECK: llvm.func @__nv_fast_powf(f32, f32) -> f32 554 // CHECK-LABEL: func @gpu_pow 555 func.func @gpu_pow(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64, f32) { 556 %result32 = math.powf %arg_f32, %arg_f32 : f32 557 // CHECK: llvm.call @__nv_powf(%{{.*}}, %{{.*}}) : (f32, f32) -> f32 558 %result64 = math.powf %arg_f64, %arg_f64 : f64 559 // CHECK: llvm.call @__nv_pow(%{{.*}}, %{{.*}}) : (f64, f64) -> f64 560 %result32Fast = math.powf %arg_f32, %arg_f32 fastmath<fast> : f32 561 // CHECK: llvm.call @__nv_fast_powf(%{{.*}}, %{{.*}}) : (f32, f32) -> f32 562 func.return %result32, %result64, %result32Fast : f32, f64, f32 563 } 564} 565 566 567 568gpu.module @test_module_27 { 569 // CHECK-LABEL: func @gpu_unroll 570 func.func @gpu_unroll(%arg0 : vector<4xf32>) -> vector<4xf32> { 571 %result = math.exp %arg0 : vector<4xf32> 572 // CHECK: %[[V0:.+]] = llvm.mlir.undef : vector<4xf32> 573 // CHECK: %[[CL:.+]] = llvm.call @__nv_expf(%{{.*}}) : (f32) -> f32 574 // CHECK: %[[V1:.+]] = llvm.insertelement %[[CL]], %[[V0]] 575 // CHECK: %[[CL:.+]] = llvm.call @__nv_expf(%{{.*}}) : (f32) -> f32 576 // CHECK: %[[V2:.+]] = llvm.insertelement %[[CL]], %[[V1]] 577 // CHECK: %[[CL:.+]] = llvm.call @__nv_expf(%{{.*}}) : (f32) -> f32 578 // CHECK: %[[V3:.+]] = llvm.insertelement %[[CL]], %[[V2]] 579 // CHECK: %[[CL:.+]] = llvm.call @__nv_expf(%{{.*}}) : (f32) -> f32 580 // CHECK: %[[V4:.+]] = llvm.insertelement %[[CL]], %[[V3]] 581 // CHECK: return %[[V4]] 582 func.return %result : vector<4xf32> 583 } 584} 585 586 587 588gpu.module @test_module_28 { 589 // CHECK-LABEL: @kernel_func 590 // CHECK: attributes 591 // CHECK: gpu.kernel 592 // CHECK: nvvm.kernel 593 gpu.func @kernel_func() kernel { 594 gpu.return 595 } 596} 597 598 599 600gpu.module @test_module_29 { 601 // CHECK-DAG: llvm.mlir.global internal constant @[[$PRINT_GLOBAL0:[A-Za-z0-9_]+]]("Hello, world\0A\00") 602 // CHECK-DAG: llvm.mlir.global internal constant @[[$PRINT_GLOBAL1:[A-Za-z0-9_]+]]("Hello: %d\0A\00") 603 // CHECK-DAG: llvm.func @vprintf(!llvm.ptr, !llvm.ptr) -> i32 604 605 // CHECK-LABEL: func @test_const_printf 606 gpu.func @test_const_printf() { 607 // CHECK-NEXT: %[[FORMATSTR:.*]] = llvm.mlir.addressof @[[$PRINT_GLOBAL0]] : !llvm.ptr 608 // CHECK-NEXT: %[[FORMATSTART:.*]] = llvm.getelementptr %[[FORMATSTR]][0, 0] : (!llvm.ptr) -> !llvm.ptr, !llvm.array<14 x i8> 609 // CHECK-NEXT: %[[O:.*]] = llvm.mlir.constant(1 : index) : i64 610 // CHECK-NEXT: %[[ALLOC:.*]] = llvm.alloca %[[O]] x !llvm.struct<()> : (i64) -> !llvm.ptr 611 // CHECK-NEXT: llvm.call @vprintf(%[[FORMATSTART]], %[[ALLOC]]) : (!llvm.ptr, !llvm.ptr) -> i32 612 gpu.printf "Hello, world\n" 613 614 // Make sure that the same global is reused. 615 // CHECK: %[[FORMATSTR2:.*]] = llvm.mlir.addressof @[[$PRINT_GLOBAL0]] : !llvm.ptr 616 // CHECK: %[[FORMATSTART2:.*]] = llvm.getelementptr %[[FORMATSTR2]][0, 0] : (!llvm.ptr) -> !llvm.ptr, !llvm.array<14 x i8> 617 // CHECK: llvm.call @vprintf(%[[FORMATSTART2]], %{{.*}}) : (!llvm.ptr, !llvm.ptr) -> i32 618 gpu.printf "Hello, world\n" 619 620 gpu.return 621 } 622 623 // CHECK-LABEL: func @test_printf 624 // CHECK: (%[[ARG0:.*]]: i32, %[[ARG1:.*]]: f32) 625 gpu.func @test_printf(%arg0: i32, %arg1: f32) { 626 // CHECK-NEXT: %[[FORMATSTR:.*]] = llvm.mlir.addressof @[[$PRINT_GLOBAL1]] : !llvm.ptr 627 // CHECK-NEXT: %[[FORMATSTART:.*]] = llvm.getelementptr %[[FORMATSTR]][0, 0] : (!llvm.ptr) -> !llvm.ptr, !llvm.array<11 x i8> 628 // CHECK-NEXT: %[[EXT:.+]] = llvm.fpext %[[ARG1]] : f32 to f64 629 // CHECK-NEXT: %[[O:.*]] = llvm.mlir.constant(1 : index) : i64 630 // CHECK-NEXT: %[[ALLOC:.*]] = llvm.alloca %[[O]] x !llvm.struct<(i32, f64)> : (i64) -> !llvm.ptr 631 // CHECK-NEXT: %[[EL0:.*]] = llvm.getelementptr %[[ALLOC]][0, 0] : (!llvm.ptr) -> !llvm.ptr, !llvm.struct<(i32, f64)> 632 // CHECK-NEXT: llvm.store %[[ARG0]], %[[EL0]] : i32, !llvm.ptr 633 // CHECK-NEXT: %[[EL1:.*]] = llvm.getelementptr %[[ALLOC]][0, 1] : (!llvm.ptr) -> !llvm.ptr, !llvm.struct<(i32, f64)> 634 // CHECK-NEXT: llvm.store %[[EXT]], %[[EL1]] : f64, !llvm.ptr 635 // CHECK-NEXT: llvm.call @vprintf(%[[FORMATSTART]], %[[ALLOC]]) : (!llvm.ptr, !llvm.ptr) -> i32 636 gpu.printf "Hello: %d\n", %arg0, %arg1 : i32, f32 637 gpu.return 638 } 639} 640 641 642 643gpu.module @test_module_30 { 644 // CHECK-LABEL: func @subgroup_reduce_add 645 gpu.func @subgroup_reduce_add(%arg0 : i32) { 646 // CHECK: nvvm.redux.sync add {{.*}} 647 %result = gpu.subgroup_reduce add %arg0 uniform {} : (i32) -> (i32) 648 gpu.return 649 } 650 // CHECK-LABEL: @subgroup_reduce_minsi 651 gpu.func @subgroup_reduce_minsi(%arg0 : i32) { 652 // CHECK: nvvm.redux.sync min {{.*}} 653 %result = gpu.subgroup_reduce minsi %arg0 uniform {} : (i32) -> (i32) 654 gpu.return 655 } 656 // CHECK-LABEL: @subgroup_reduce_maxsi 657 gpu.func @subgroup_reduce_maxsi(%arg0 : i32) { 658 // CHECK: nvvm.redux.sync max {{.*}} 659 %result = gpu.subgroup_reduce maxsi %arg0 uniform {} : (i32) -> (i32) 660 gpu.return 661 } 662 // CHECK-LABEL: func @subgroup_reduce_and 663 gpu.func @subgroup_reduce_and(%arg0 : i32) { 664 // CHECK: nvvm.redux.sync and {{.*}} 665 %result = gpu.subgroup_reduce and %arg0 uniform {} : (i32) -> (i32) 666 gpu.return 667 } 668 // CHECK-LABEL: @subgroup_reduce_or 669 gpu.func @subgroup_reduce_or(%arg0 : i32) { 670 // CHECK: nvvm.redux.sync or {{.*}} 671 %result = gpu.subgroup_reduce or %arg0 uniform {} : (i32) -> (i32) 672 gpu.return 673 } 674 // CHECK-LABEL: @subgroup_reduce_xor 675 gpu.func @subgroup_reduce_xor(%arg0 : i32) { 676 // CHECK: nvvm.redux.sync xor {{.*}} 677 %result = gpu.subgroup_reduce xor %arg0 uniform {} : (i32) -> (i32) 678 gpu.return 679 } 680} 681 682gpu.module @test_module_31 { 683 // CHECK: llvm.func @__nv_fmodf(f32, f32) -> f32 684 // CHECK: llvm.func @__nv_fmod(f64, f64) -> f64 685 // CHECK-LABEL: func @gpu_fmod 686 func.func @gpu_fmod(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) { 687 %result32 = arith.remf %arg_f32, %arg_f32 : f32 688 // CHECK: llvm.call @__nv_fmodf(%{{.*}}, %{{.*}}) : (f32, f32) -> f32 689 %result64 = arith.remf %arg_f64, %arg_f64 : f64 690 // CHECK: llvm.call @__nv_fmod(%{{.*}}, %{{.*}}) : (f64, f64) -> f64 691 func.return %result32, %result64 : f32, f64 692 } 693} 694 695gpu.module @test_module_32 { 696 // CHECK: llvm.func @__nv_erff(f32) -> f32 697 // CHECK: llvm.func @__nv_erf(f64) -> f64 698 // CHECK-LABEL: func @gpu_erf 699 func.func @gpu_erf(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) { 700 %result32 = math.erf %arg_f32 : f32 701 // CHECK: llvm.call @__nv_erff(%{{.*}}) : (f32) -> f32 702 %result64 = math.erf %arg_f64 : f64 703 // CHECK: llvm.call @__nv_erf(%{{.*}}) : (f64) -> f64 704 func.return %result32, %result64 : f32, f64 705 } 706} 707 708gpu.module @test_module_33 { 709// CHECK-LABEL: func @kernel_with_block_size( 710// CHECK: attributes {gpu.kernel, gpu.known_block_size = array<i32: 32, 4, 2>, nvvm.kernel, nvvm.maxntid = array<i32: 32, 4, 2>} 711 gpu.func @kernel_with_block_size(%arg0: !llvm.ptr) kernel attributes {known_block_size = array<i32: 32, 4, 2>} { 712 // CHECK: = nvvm.read.ptx.sreg.tid.x range <i32, 0, 32> : i32 713 %0 = gpu.thread_id x 714 // CHECK: = nvvm.read.ptx.sreg.tid.y range <i32, 0, 4> : i32 715 %1 = gpu.thread_id y 716 // CHECK: = nvvm.read.ptx.sreg.tid.z range <i32, 0, 2> : i32 717 %2 = gpu.thread_id z 718 719 // Fake usage to prevent dead code elimination 720 %3 = arith.addi %0, %1 : index 721 %4 = arith.addi %3, %2 : index 722 %5 = arith.index_cast %4 : index to i64 723 llvm.store %5, %arg0 : i64, !llvm.ptr 724 gpu.return 725 } 726} 727 728 729gpu.module @test_module_34 { 730 // CHECK-LABEL: llvm.func @memref_signature( 731 // CHECK-SAME: %{{.*}}: !llvm.ptr, %{{.*}}: !llvm.ptr, %{{.*}}: i64, %{{.*}}: i64, %{{.*}}: i64, %{{.*}}: f32) -> !llvm.struct<(struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)>, f32)> 732 // CHECK: llvm.mlir.undef 733 // CHECK: llvm.insertvalue 734 // CHECK: llvm.insertvalue 735 // CHECK: llvm.insertvalue 736 // CHECK: llvm.insertvalue 737 // CHECK: llvm.insertvalue 738 // CHECK: llvm.mlir.undef 739 // CHECK: llvm.insertvalue 740 // CHECK: llvm.insertvalue 741 // CHECK: llvm.return 742 743 // CHECK-BARE-LABEL: llvm.func @memref_signature( 744 // CHECK-BARE-SAME: %{{.*}}: !llvm.ptr, %{{.*}}: f32) -> !llvm.struct<(ptr, f32)> 745 gpu.func @memref_signature(%m: memref<2xf32>, %f: f32) -> (memref<2xf32>, f32) { 746 gpu.return %m, %f : memref<2xf32>, f32 747 } 748} 749 750gpu.module @test_module_35 { 751 // CHECK: llvm.func @__nv_acosf(f32) -> f32 752 // CHECK: llvm.func @__nv_acos(f64) -> f64 753 // CHECK-LABEL: func @gpu_acos 754 func.func @gpu_acos(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) { 755 %result32 = math.acos %arg_f32 : f32 756 // CHECK: llvm.call @__nv_acosf(%{{.*}}) : (f32) -> f32 757 %result64 = math.acos %arg_f64 : f64 758 // CHECK: llvm.call @__nv_acos(%{{.*}}) : (f64) -> f64 759 func.return %result32, %result64 : f32, f64 760 } 761} 762 763gpu.module @test_module_36 { 764 // CHECK: llvm.func @__nv_acoshf(f32) -> f32 765 // CHECK: llvm.func @__nv_acosh(f64) -> f64 766 // CHECK-LABEL: func @gpu_acosh 767 func.func @gpu_acosh(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) { 768 %result32 = math.acosh %arg_f32 : f32 769 // CHECK: llvm.call @__nv_acoshf(%{{.*}}) : (f32) -> f32 770 %result64 = math.acosh %arg_f64 : f64 771 // CHECK: llvm.call @__nv_acosh(%{{.*}}) : (f64) -> f64 772 func.return %result32, %result64 : f32, f64 773 } 774} 775 776gpu.module @test_module_37 { 777 // CHECK: llvm.func @__nv_asinf(f32) -> f32 778 // CHECK: llvm.func @__nv_asin(f64) -> f64 779 // CHECK-LABEL: func @gpu_asin 780 func.func @gpu_asin(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) { 781 %result32 = math.asin %arg_f32 : f32 782 // CHECK: llvm.call @__nv_asinf(%{{.*}}) : (f32) -> f32 783 %result64 = math.asin %arg_f64 : f64 784 // CHECK: llvm.call @__nv_asin(%{{.*}}) : (f64) -> f64 785 func.return %result32, %result64 : f32, f64 786 } 787} 788 789gpu.module @test_module_38 { 790 // CHECK: llvm.func @__nv_asinhf(f32) -> f32 791 // CHECK: llvm.func @__nv_asinh(f64) -> f64 792 // CHECK-LABEL: func @gpu_asinh 793 func.func @gpu_asinh(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) { 794 %result32 = math.asinh %arg_f32 : f32 795 // CHECK: llvm.call @__nv_asinhf(%{{.*}}) : (f32) -> f32 796 %result64 = math.asinh %arg_f64 : f64 797 // CHECK: llvm.call @__nv_asinh(%{{.*}}) : (f64) -> f64 798 func.return %result32, %result64 : f32, f64 799 } 800} 801 802gpu.module @test_module_39 { 803 // CHECK: llvm.func @__nv_atanhf(f32) -> f32 804 // CHECK: llvm.func @__nv_atanh(f64) -> f64 805 // CHECK-LABEL: func @gpu_atanh 806 func.func @gpu_atanh(%arg_f16 : f16, %arg_f32 : f32, %arg_f64 : f64) 807 -> (f16, f32, f64) { 808 %result16 = math.atanh %arg_f16 : f16 809 // CHECK: llvm.fpext %{{.*}} : f16 to f32 810 // CHECK-NEXT: llvm.call @__nv_atanhf(%{{.*}}) : (f32) -> f32 811 // CHECK-NEXT: llvm.fptrunc %{{.*}} : f32 to f16 812 %result32 = math.atanh %arg_f32 : f32 813 // CHECK: llvm.call @__nv_atanhf(%{{.*}}) : (f32) -> f32 814 %result64 = math.atanh %arg_f64 : f64 815 // CHECK: llvm.call @__nv_atanh(%{{.*}}) : (f64) -> f64 816 func.return %result16, %result32, %result64 : f16, f32, f64 817 } 818} 819 820gpu.module @test_module_40 { 821 // CHECK: llvm.func @__nv_copysignf(f32, f32) -> f32 822 // CHECK: llvm.func @__nv_copysign(f64, f64) -> f64 823 // CHECK-LABEL: func @gpu_copysign 824 func.func @gpu_copysign(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) { 825 %result32 = math.copysign %arg_f32, %arg_f32 : f32 826 // CHECK: llvm.call @__nv_copysignf(%{{.*}}, %{{.*}}) : (f32, f32) -> f32 827 %result64 = math.copysign %arg_f64, %arg_f64 : f64 828 // CHECK: llvm.call @__nv_copysign(%{{.*}}, %{{.*}}) : (f64, f64) -> f64 829 func.return %result32, %result64 : f32, f64 830 } 831} 832 833gpu.module @test_module_41 { 834 // CHECK: llvm.func @__nv_coshf(f32) -> f32 835 // CHECK: llvm.func @__nv_cosh(f64) -> f64 836 // CHECK-LABEL: func @gpu_cosh 837 func.func @gpu_cosh(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) { 838 %result32 = math.cosh %arg_f32 : f32 839 // CHECK: llvm.call @__nv_coshf(%{{.*}}) : (f32) -> f32 840 %result64 = math.cosh %arg_f64 : f64 841 // CHECK: llvm.call @__nv_cosh(%{{.*}}) : (f64) -> f64 842 func.return %result32, %result64 : f32, f64 843 } 844} 845 846gpu.module @test_module_42 { 847 // CHECK: llvm.func @__nv_fmaf(f32, f32, f32) -> f32 848 // CHECK: llvm.func @__nv_fma(f64, f64, f64) -> f64 849 // CHECK-LABEL: func @gpu_fma 850 func.func @gpu_fma(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) { 851 %result32 = math.fma %arg_f32, %arg_f32, %arg_f32 : f32 852 // CHECK: llvm.call @__nv_fmaf(%{{.*}}, %{{.*}}, %{{.*}}) : (f32, f32, f32) -> f32 853 %result64 = math.fma %arg_f64, %arg_f64, %arg_f64 : f64 854 // CHECK: llvm.call @__nv_fma(%{{.*}}, %{{.*}}, %{{.*}}) : (f64, f64, f64) -> f64 855 func.return %result32, %result64 : f32, f64 856 } 857} 858 859gpu.module @test_module_43 { 860 // CHECK: llvm.func @__nv_roundf(f32) -> f32 861 // CHECK: llvm.func @__nv_round(f64) -> f64 862 // CHECK-LABEL: func @gpu_round 863 func.func @gpu_round(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) { 864 %result32 = math.round %arg_f32 : f32 865 // CHECK: llvm.call @__nv_roundf(%{{.*}}) : (f32) -> f32 866 %result64 = math.round %arg_f64 : f64 867 // CHECK: llvm.call @__nv_round(%{{.*}}) : (f64) -> f64 868 func.return %result32, %result64 : f32, f64 869 } 870} 871 872gpu.module @test_module_44 { 873 // CHECK: llvm.func @__nv_rintf(f32) -> f32 874 // CHECK: llvm.func @__nv_rint(f64) -> f64 875 // CHECK-LABEL: func @gpu_roundeven 876 func.func @gpu_roundeven(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) { 877 %result32 = math.roundeven %arg_f32 : f32 878 // CHECK: llvm.call @__nv_rintf(%{{.*}}) : (f32) -> f32 879 %result64 = math.roundeven %arg_f64 : f64 880 // CHECK: llvm.call @__nv_rint(%{{.*}}) : (f64) -> f64 881 func.return %result32, %result64 : f32, f64 882 } 883} 884 885gpu.module @test_module_45 { 886 // CHECK: llvm.func @__nv_sinhf(f32) -> f32 887 // CHECK: llvm.func @__nv_sinh(f64) -> f64 888 // CHECK-LABEL: func @gpu_sinh 889 func.func @gpu_sinh(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) { 890 %result32 = math.sinh %arg_f32 : f32 891 // CHECK: llvm.call @__nv_sinhf(%{{.*}}) : (f32) -> f32 892 %result64 = math.sinh %arg_f64 : f64 893 // CHECK: llvm.call @__nv_sinh(%{{.*}}) : (f64) -> f64 894 func.return %result32, %result64 : f32, f64 895 } 896} 897 898gpu.module @test_module_46 { 899 // CHECK: llvm.func @__nv_coshf(f32) -> f32 900 // CHECK: llvm.func @__nv_cosh(f64) -> f64 901 // CHECK-LABEL: func @gpu_cosh_with_fastmath 902 func.func @gpu_cosh_with_fastmath(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) { 903 %result32 = math.cosh %arg_f32 fastmath<fast> : f32 904 // CHECK: llvm.call @__nv_coshf(%{{.*}}) : (f32) -> f32 905 %result64 = math.cosh %arg_f64 fastmath<afn> : f64 906 // CHECK: llvm.call @__nv_cosh(%{{.*}}) : (f64) -> f64 907 func.return %result32, %result64 : f32, f64 908 } 909} 910 911gpu.module @test_module_47 { 912 // CHECK: llvm.func @__nv_sinhf(f32) -> f32 913 // CHECK: llvm.func @__nv_sinh(f64) -> f64 914 // CHECK-LABEL: func @gpu_sinh_with_fastmath 915 func.func @gpu_sinh_with_fastmath(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) { 916 %result32 = math.sinh %arg_f32 fastmath<contract> : f32 917 // CHECK: llvm.call @__nv_sinhf(%{{.*}}) : (f32) -> f32 918 %result64 = math.sinh %arg_f64 fastmath<none> : f64 919 // CHECK: llvm.call @__nv_sinh(%{{.*}}) : (f64) -> f64 920 func.return %result32, %result64 : f32, f64 921 } 922} 923 924gpu.module @test_module_48 { 925 // CHECK: llvm.func @__nv_expf(f32) -> f32 926 // CHECK: llvm.func @__nv_exp(f64) -> f64 927 // CHECK-LABEL: func @gpu_exp_with_fastmath 928 func.func @gpu_exp_with_fastmath(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64, f32) { 929 %result32 = math.exp %arg_f32 fastmath<reassoc>: f32 930 // CHECK: llvm.call @__nv_expf(%{{.*}}) : (f32) -> f32 931 %result64 = math.exp %arg_f64 fastmath<contract>: f64 932 // CHECK: llvm.call @__nv_exp(%{{.*}}) : (f64) -> f64 933 %result32Fast = math.exp %arg_f32 fastmath<ninf> : f32 934 // CHECK: llvm.call @__nv_expf(%{{.*}}) : (f32) -> f32 935 func.return %result32, %result64, %result32Fast : f32, f64, f32 936 } 937} 938 939gpu.module @test_module_49 { 940// CHECK-LABEL: func @explicit_id_bounds() 941 func.func @explicit_id_bounds() -> (index, index, index) { 942 // CHECK: = nvvm.read.ptx.sreg.tid.x range <i32, 0, 32> : i32 943 %0 = gpu.thread_id x upper_bound 32 944 // CHECK: = nvvm.read.ptx.sreg.ntid.x range <i32, 1, 33> : i32 945 %1 = gpu.block_dim x upper_bound 32 946 // CHECK: = nvvm.read.ptx.sreg.laneid range <i32, 0, 16> : i32 947 %2 = gpu.lane_id upper_bound 16 948 949 return %0, %1, %2 : index, index, index 950 } 951} 952 953gpu.module @test_module_50 { 954// CHECK-LABEL: func @kernel_with_grid_size( 955 gpu.func @kernel_with_grid_size(%arg0: !llvm.ptr) kernel attributes {known_grid_size = array<i32: 32, 4, 2>} { 956 // CHECK: = nvvm.read.ptx.sreg.ctaid.x range <i32, 0, 32> : i32 957 %0 = gpu.block_id x 958 // CHECK: = nvvm.read.ptx.sreg.ctaid.y range <i32, 0, 4> : i32 959 %1 = gpu.block_id y 960 // CHECK: = nvvm.read.ptx.sreg.ctaid.z range <i32, 0, 2> : i32 961 %2 = gpu.block_id z 962 963 // Fake usage to prevent dead code elimination 964 %3 = arith.addi %0, %1 : index 965 %4 = arith.addi %3, %2 : index 966 %5 = arith.index_cast %4 : index to i64 967 llvm.store %5, %arg0 : i64, !llvm.ptr 968 gpu.return 969 } 970} 971 972// CHECK-LABEL: gpu.module @test_module_51 973// CHECK: llvm.mlir.global internal constant @[[func_name:.*]]("(unknown)\00") {addr_space = 0 : i32} 974// CHECK: llvm.mlir.global internal constant @[[file_name:.*]]("{{.*}}gpu-to-nvvm.mlir{{.*}}") {addr_space = 0 : i32} 975// CHECK: llvm.mlir.global internal constant @[[message:.*]]("assert message\00") {addr_space = 0 : i32} 976// CHECK: llvm.func @__assertfail(!llvm.ptr, !llvm.ptr, i32, !llvm.ptr, i64) attributes {passthrough = ["noreturn"]} 977// CHECK: llvm.func @test_assert(%[[cond:.*]]: i1) attributes {gpu.kernel, nvvm.kernel} { 978// CHECK: llvm.cond_br %[[cond]], ^[[after_block:.*]], ^[[assert_block:.*]] 979// CHECK: ^[[assert_block]]: 980// CHECK: %[[message_ptr:.*]] = llvm.mlir.addressof @[[message]] : !llvm.ptr 981// CHECK: %[[message_start:.*]] = llvm.getelementptr %[[message_ptr]][0, 0] : (!llvm.ptr) -> !llvm.ptr, !llvm.array<15 x i8> 982// CHECK: %[[file_ptr:.*]] = llvm.mlir.addressof @[[file_name]] : !llvm.ptr 983// CHECK: %[[file_start:.*]] = llvm.getelementptr %[[file_ptr]][0, 0] : (!llvm.ptr) -> !llvm.ptr, !llvm.array<{{.*}} x i8> 984// CHECK: %[[func_ptr:.*]] = llvm.mlir.addressof @[[func_name]] : !llvm.ptr 985// CHECK: %[[func_start:.*]] = llvm.getelementptr %[[func_ptr]][0, 0] : (!llvm.ptr) -> !llvm.ptr, !llvm.array<{{.*}} x i8> 986// CHECK: %[[line_num:.*]] = llvm.mlir.constant({{.*}} : i32) : i32 987// CHECK: %[[ptr:.*]] = llvm.mlir.constant(1 : i64) : i64 988// CHECK: llvm.call @__assertfail(%[[message_start]], %[[file_start]], %[[line_num]], %[[func_start]], %[[ptr]]) : (!llvm.ptr, !llvm.ptr, i32, !llvm.ptr, i64) -> () 989// CHECK: llvm.br ^[[after_block]] 990// CHECK: ^[[after_block]]: 991// CHECK: llvm.return 992// CHECK: } 993 994gpu.module @test_module_51 { 995 gpu.func @test_assert(%arg0: i1) kernel { 996 cf.assert %arg0, "assert message" 997 gpu.return 998 } 999} 1000 1001module attributes {transform.with_named_sequence} { 1002 transform.named_sequence @__transform_main(%toplevel_module: !transform.any_op {transform.readonly}) { 1003 %gpu_module = transform.structured.match ops{["gpu.module"]} in %toplevel_module 1004 : (!transform.any_op) -> !transform.any_op 1005 1006 transform.apply_patterns to %gpu_module { 1007 transform.apply_patterns.gpu.gpu_rewrite_patterns 1008 } : !transform.any_op 1009 1010 transform.apply_conversion_patterns to %gpu_module { 1011 transform.apply_conversion_patterns.dialect_to_llvm "arith" 1012 transform.apply_conversion_patterns.dialect_to_llvm "cf" 1013 transform.apply_conversion_patterns.vector.vector_to_llvm 1014 transform.apply_conversion_patterns.func.func_to_llvm 1015 transform.apply_conversion_patterns.dialect_to_llvm "memref" 1016 transform.apply_conversion_patterns.gpu.gpu_to_nvvm 1017 transform.apply_conversion_patterns.gpu.gpu_wmma_to_nvvm 1018 transform.apply_conversion_patterns.gpu.gpu_subgroup_reduce_to_nvvm 1019 transform.apply_conversion_patterns.nvgpu.nvgpu_to_nvvm 1020 } with type_converter { 1021 transform.apply_conversion_patterns.memref.memref_to_llvm_type_converter 1022 {index_bitwidth = 64, 1023 use_bare_ptr_call_conv = false} 1024 } { 1025 legal_dialects = ["llvm", "memref", "nvvm", "test"], 1026 legal_ops = ["func.func", "gpu.module", "gpu.yield"], 1027 illegal_dialects = ["gpu"], 1028 illegal_ops = ["llvm.copysign", "llvm.cos", "llvm.exp", "llvm.exp2", "llvm.fabs", "llvm.fceil", 1029 "llvm.ffloor", "llvm.fma", "llvm.frem", "llvm.log", "llvm.log10", "llvm.log2", "llvm.pow", 1030 "llvm.roundeven", "llvm.round", "llvm.sin", "llvm.sqrt"], 1031 partial_conversion 1032 } : !transform.any_op 1033 transform.yield 1034 } 1035} 1036 1037 1038gpu.module @test_module_52 { 1039 // CHECK: llvm.func @__nv_abs(i32) -> i32 1040 // CHECK-LABEL: func @gpu_abs 1041 func.func @gpu_abs(%arg_i32 : i32) -> (i32) { 1042 %result32 = math.absi %arg_i32 : i32 1043 // CHECK: llvm.call @__nv_abs(%{{.*}}) : (i32) -> i32 1044 func.return %result32 : i32 1045 } 1046} 1047 1048gpu.module @test_module_53 { 1049 // CHECK: llvm.func @__nv_powif(f32, i32) -> f32 1050 // CHECK: llvm.func @__nv_powi(f64, i32) -> f64 1051 // CHECK-LABEL: func @gpu_powi 1052 func.func @gpu_powi(%arg_f32 : f32, %arg_f64 : f64, %arg_i32 : i32) -> (f32, f64) { 1053 %result32 = math.fpowi %arg_f32, %arg_i32 : f32, i32 1054 // CHECK: llvm.call @__nv_powif(%{{.*}}, %{{.*}}) : (f32, i32) -> f32 1055 %result64 = math.fpowi %arg_f64, %arg_i32 : f64, i32 1056 // CHECK: llvm.call @__nv_powi(%{{.*}}, %{{.*}}) : (f64, i32) -> f64 1057 func.return %result32, %result64 : f32, f64 1058 } 1059} 1060