mlir/test/Conversion/GPUToNVVM/gpu-to-nvvm.mlir

   1 // RUN: mlir-opt %s -convert-gpu-to-nvvm -split-input-file | FileCheck %s
   2 // RUN: mlir-opt %s -convert-gpu-to-nvvm='index-bitwidth=32' -split-input-file | FileCheck --check-prefix=CHECK32 %s
   3
   4 gpu.module @test_module {
   5   // CHECK-LABEL: func @gpu_index_ops()
   6   // CHECK32-LABEL: func @gpu_index_ops()
   7   builtin.func @gpu_index_ops()
   8       -> (index, index, index, index, index, index,
   9           index, index, index, index, index, index) {
  10     // CHECK32-NOT: = llvm.sext %{{.*}} : i32 to i64
  11
  12     // CHECK: = nvvm.read.ptx.sreg.tid.x : i32
  13     // CHECK: = llvm.sext %{{.*}} : i32 to i64
  14     %tIdX = "gpu.thread_id"() {dimension = "x"} : () -> (index)
  15     // CHECK: = nvvm.read.ptx.sreg.tid.y : i32
  16     // CHECK: = llvm.sext %{{.*}} : i32 to i64
  17     %tIdY = "gpu.thread_id"() {dimension = "y"} : () -> (index)
  18     // CHECK: = nvvm.read.ptx.sreg.tid.z : i32
  19     // CHECK: = llvm.sext %{{.*}} : i32 to i64
  20     %tIdZ = "gpu.thread_id"() {dimension = "z"} : () -> (index)
  21
  22     // CHECK: = nvvm.read.ptx.sreg.ntid.x : i32
  23     // CHECK: = llvm.sext %{{.*}} : i32 to i64
  24     %bDimX = "gpu.block_dim"() {dimension = "x"} : () -> (index)
  25     // CHECK: = nvvm.read.ptx.sreg.ntid.y : i32
  26     // CHECK: = llvm.sext %{{.*}} : i32 to i64
  27     %bDimY = "gpu.block_dim"() {dimension = "y"} : () -> (index)
  28     // CHECK: = nvvm.read.ptx.sreg.ntid.z : i32
  29     // CHECK: = llvm.sext %{{.*}} : i32 to i64
  30     %bDimZ = "gpu.block_dim"() {dimension = "z"} : () -> (index)
  31
  32     // CHECK: = nvvm.read.ptx.sreg.ctaid.x : i32
  33     // CHECK: = llvm.sext %{{.*}} : i32 to i64
  34     %bIdX = "gpu.block_id"() {dimension = "x"} : () -> (index)
  35     // CHECK: = nvvm.read.ptx.sreg.ctaid.y : i32
  36     // CHECK: = llvm.sext %{{.*}} : i32 to i64
  37     %bIdY = "gpu.block_id"() {dimension = "y"} : () -> (index)
  38     // CHECK: = nvvm.read.ptx.sreg.ctaid.z : i32
  39     // CHECK: = llvm.sext %{{.*}} : i32 to i64
  40     %bIdZ = "gpu.block_id"() {dimension = "z"} : () -> (index)
  41
  42     // CHECK: = nvvm.read.ptx.sreg.nctaid.x : i32
  43     // CHECK: = llvm.sext %{{.*}} : i32 to i64
  44     %gDimX = "gpu.grid_dim"() {dimension = "x"} : () -> (index)
  45     // CHECK: = nvvm.read.ptx.sreg.nctaid.y : i32
  46     // CHECK: = llvm.sext %{{.*}} : i32 to i64
  47     %gDimY = "gpu.grid_dim"() {dimension = "y"} : () -> (index)
  48     // CHECK: = nvvm.read.ptx.sreg.nctaid.z : i32
  49     // CHECK: = llvm.sext %{{.*}} : i32 to i64
  50     %gDimZ = "gpu.grid_dim"() {dimension = "z"} : () -> (index)
  51
  52     std.return %tIdX, %tIdY, %tIdZ, %bDimX, %bDimY, %bDimZ,
  53                %bIdX, %bIdY, %bIdZ, %gDimX, %gDimY, %gDimZ
  54         : index, index, index, index, index, index,
  55           index, index, index, index, index, index
  56   }
  57 }
  58
  59 // -----
  60
  61 gpu.module @test_module {
  62   // CHECK-LABEL: func @gpu_index_comp
  63   // CHECK32-LABEL: func @gpu_index_comp
  64   builtin.func @gpu_index_comp(%idx : index) -> index {
  65     // CHECK: = llvm.add %{{.*}}, %{{.*}} : i64
  66     // CHECK32: = llvm.add %{{.*}}, %{{.*}} : i32
  67     %0 = arith.addi %idx, %idx : index
  68     // CHECK: llvm.return %{{.*}} : i64
  69     // CHECK32: llvm.return %{{.*}} : i32
  70     std.return %0 : index
  71   }
  72 }
  73
  74 // -----
  75
  76 gpu.module @test_module {
  77   // CHECK-LABEL: func @gpu_all_reduce_op()
  78   gpu.func @gpu_all_reduce_op() {
  79     %arg0 = arith.constant 1.0 : f32
  80     // TODO: Check full IR expansion once lowering has settled.
  81     // CHECK: nvvm.shfl.sync "bfly" {{.*}}
  82     // CHECK: nvvm.barrier0
  83     // CHECK: llvm.fadd
  84     %result = "gpu.all_reduce"(%arg0) ({}) {op = "add"} : (f32) -> (f32)
  85
  86     gpu.return
  87   }
  88 }
  89
  90 // -----
  91
  92 gpu.module @test_module {
  93   // CHECK-LABEL: func @gpu_all_reduce_region()
  94   gpu.func @gpu_all_reduce_region() {
  95     %arg0 = arith.constant 1 : i32
  96     // TODO: Check full IR expansion once lowering has settled.
  97     // CHECK: nvvm.shfl.sync "bfly" {{.*}}
  98     // CHECK: nvvm.barrier0
  99     %result = "gpu.all_reduce"(%arg0) ({
 100     ^bb(%lhs : i32, %rhs : i32):
 101       %xor = arith.xori %lhs, %rhs : i32
 102       "gpu.yield"(%xor) : (i32) -> ()
 103     }) : (i32) -> (i32)
 104     gpu.return
 105   }
 106 }
 107
 108 // -----
 109
 110 gpu.module @test_module {
 111   // CHECK-LABEL: func @gpu_shuffle()
 112   builtin.func @gpu_shuffle() -> (f32, f32, f32, f32) {
 113     // CHECK: %[[#VALUE:]] = llvm.mlir.constant(1.000000e+00 : f32) : f32
 114     %arg0 = arith.constant 1.0 : f32
 115     // CHECK: %[[#OFFSET:]] = llvm.mlir.constant(4 : i32) : i32
 116     %arg1 = arith.constant 4 : i32
 117     // CHECK: %[[#WIDTH:]] = llvm.mlir.constant(23 : i32) : i32
 118     %arg2 = arith.constant 23 : i32
 119     // CHECK: %[[#ONE:]] = llvm.mlir.constant(1 : i32) : i32
 120     // CHECK: %[[#SHL:]] = llvm.shl %[[#ONE]], %[[#WIDTH]] : i32
 121     // CHECK: %[[#MASK:]] = llvm.sub %[[#SHL]], %[[#ONE]] : i32
 122     // CHECK: %[[#CLAMP:]] = llvm.sub %[[#WIDTH]], %[[#ONE]] : i32
 123     // CHECK: %[[#SHFL:]] = nvvm.shfl.sync "bfly" %[[#MASK]], %[[#VALUE]], %[[#OFFSET]], %[[#CLAMP]] {return_value_and_is_valid} : f32 -> !llvm.struct<(f32, i1)>
 124     // CHECK: llvm.extractvalue %[[#SHFL]][0 : index] : !llvm.struct<(f32, i1)>
 125     // CHECK: llvm.extractvalue %[[#SHFL]][1 : index] : !llvm.struct<(f32, i1)>
 126     %shfl, %pred = "gpu.shuffle"(%arg0, %arg1, %arg2) { mode = "xor" } : (f32, i32, i32) -> (f32, i1)
 127     // CHECK: nvvm.shfl.sync "up" {{.*}} {return_value_and_is_valid} : f32 -> !llvm.struct<(f32, i1)>
 128     %shflu, %predu = "gpu.shuffle"(%arg0, %arg1, %arg2) { mode = "up" } : (f32, i32, i32) -> (f32, i1)
 129     // CHECK: nvvm.shfl.sync "down" {{.*}} {return_value_and_is_valid} : f32 -> !llvm.struct<(f32, i1)>
 130     %shfld, %predd = "gpu.shuffle"(%arg0, %arg1, %arg2) { mode = "down" } : (f32, i32, i32) -> (f32, i1)
 131     // CHECK: nvvm.shfl.sync "idx" {{.*}} {return_value_and_is_valid} : f32 -> !llvm.struct<(f32, i1)>
 132     %shfli, %predi = "gpu.shuffle"(%arg0, %arg1, %arg2) { mode = "idx" } : (f32, i32, i32) -> (f32, i1)
 133
 134     std.return %shfl, %shflu, %shfld, %shfli : f32, f32,f32, f32
 135   }
 136 }
 137
 138 // -----
 139
 140 gpu.module @test_module {
 141   // CHECK-LABEL: func @gpu_sync()
 142   builtin.func @gpu_sync() {
 143     // CHECK: nvvm.barrier0
 144     gpu.barrier
 145     std.return
 146   }
 147 }
 148
 149 // -----
 150
 151 gpu.module @test_module {
 152   // CHECK: llvm.func @__nv_fabsf(f32) -> f32
 153   // CHECK: llvm.func @__nv_fabs(f64) -> f64
 154   // CHECK-LABEL: func @gpu_fabs
 155   builtin.func @gpu_fabs(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
 156     %result32 = math.abs %arg_f32 : f32
 157     // CHECK: llvm.call @__nv_fabsf(%{{.*}}) : (f32) -> f32
 158     %result64 = math.abs %arg_f64 : f64
 159     // CHECK: llvm.call @__nv_fabs(%{{.*}}) : (f64) -> f64
 160     std.return %result32, %result64 : f32, f64
 161   }
 162 }
 163
 164 // -----
 165
 166 gpu.module @test_module {
 167   // CHECK: llvm.func @__nv_ceilf(f32) -> f32
 168   // CHECK: llvm.func @__nv_ceil(f64) -> f64
 169   // CHECK-LABEL: func @gpu_ceil
 170   builtin.func @gpu_ceil(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
 171     %result32 = math.ceil %arg_f32 : f32
 172     // CHECK: llvm.call @__nv_ceilf(%{{.*}}) : (f32) -> f32
 173     %result64 = math.ceil %arg_f64 : f64
 174     // CHECK: llvm.call @__nv_ceil(%{{.*}}) : (f64) -> f64
 175     std.return %result32, %result64 : f32, f64
 176   }
 177 }
 178
 179 // -----
 180
 181 gpu.module @test_module {
 182   // CHECK: llvm.func @__nv_floorf(f32) -> f32
 183   // CHECK: llvm.func @__nv_floor(f64) -> f64
 184   // CHECK-LABEL: func @gpu_floor
 185   builtin.func @gpu_floor(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
 186     %result32 = math.floor %arg_f32 : f32
 187     // CHECK: llvm.call @__nv_floorf(%{{.*}}) : (f32) -> f32
 188     %result64 = math.floor %arg_f64 : f64
 189     // CHECK: llvm.call @__nv_floor(%{{.*}}) : (f64) -> f64
 190     std.return %result32, %result64 : f32, f64
 191   }
 192 }
 193
 194 // -----
 195
 196 gpu.module @test_module {
 197   // CHECK: llvm.func @__nv_cosf(f32) -> f32
 198   // CHECK: llvm.func @__nv_cos(f64) -> f64
 199   // CHECK-LABEL: func @gpu_cos
 200   builtin.func @gpu_cos(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
 201     %result32 = math.cos %arg_f32 : f32
 202     // CHECK: llvm.call @__nv_cosf(%{{.*}}) : (f32) -> f32
 203     %result64 = math.cos %arg_f64 : f64
 204     // CHECK: llvm.call @__nv_cos(%{{.*}}) : (f64) -> f64
 205     std.return %result32, %result64 : f32, f64
 206   }
 207 }
 208
 209 // -----
 210 gpu.module @test_module {
 211   // CHECK: llvm.func @__nv_expf(f32) -> f32
 212   // CHECK: llvm.func @__nv_exp(f64) -> f64
 213   // CHECK-LABEL: func @gpu_exp
 214   builtin.func @gpu_exp(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
 215     %result32 = math.exp %arg_f32 : f32
 216     // CHECK: llvm.call @__nv_expf(%{{.*}}) : (f32) -> f32
 217     %result64 = math.exp %arg_f64 : f64
 218     // CHECK: llvm.call @__nv_exp(%{{.*}}) : (f64) -> f64
 219     std.return %result32, %result64 : f32, f64
 220   }
 221 }
 222
 223 // -----
 224 gpu.module @test_module {
 225   // CHECK: llvm.func @__nv_exp2f(f32) -> f32
 226   // CHECK: llvm.func @__nv_exp2(f64) -> f64
 227   // CHECK-LABEL: func @gpu_exp2
 228   builtin.func @gpu_exp2(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
 229     %result32 = math.exp2 %arg_f32 : f32
 230     // CHECK: llvm.call @__nv_exp2f(%{{.*}}) : (f32) -> f32
 231     %result64 = math.exp2 %arg_f64 : f64
 232     // CHECK: llvm.call @__nv_exp2(%{{.*}}) : (f64) -> f64
 233     std.return %result32, %result64 : f32, f64
 234   }
 235 }
 236
 237 // -----
 238
 239 gpu.module @test_module {
 240   // CHECK: llvm.func @__nv_logf(f32) -> f32
 241   // CHECK: llvm.func @__nv_log(f64) -> f64
 242   // CHECK-LABEL: func @gpu_log
 243   builtin.func @gpu_log(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
 244     %result32 = math.log %arg_f32 : f32
 245     // CHECK: llvm.call @__nv_logf(%{{.*}}) : (f32) -> f32
 246     %result64 = math.log %arg_f64 : f64
 247     // CHECK: llvm.call @__nv_log(%{{.*}}) : (f64) -> f64
 248     std.return %result32, %result64 : f32, f64
 249   }
 250 }
 251
 252 // -----
 253
 254 gpu.module @test_module {
 255   // CHECK: llvm.func @__nv_log10f(f32) -> f32
 256   // CHECK: llvm.func @__nv_log10(f64) -> f64
 257   // CHECK-LABEL: func @gpu_log10
 258   builtin.func @gpu_log10(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
 259     %result32 = math.log10 %arg_f32 : f32
 260     // CHECK: llvm.call @__nv_log10f(%{{.*}}) : (f32) -> f32
 261     %result64 = math.log10 %arg_f64 : f64
 262     // CHECK: llvm.call @__nv_log10(%{{.*}}) : (f64) -> f64
 263     std.return %result32, %result64 : f32, f64
 264   }
 265 }
 266
 267 // -----
 268
 269 gpu.module @test_module {
 270   // CHECK: llvm.func @__nv_log1pf(f32) -> f32
 271   // CHECK: llvm.func @__nv_log1p(f64) -> f64
 272   // CHECK-LABEL: func @gpu_log1p
 273   builtin.func @gpu_log1p(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
 274     %result32 = math.log1p %arg_f32 : f32
 275     // CHECK: llvm.call @__nv_log1pf(%{{.*}}) : (f32) -> f32
 276     %result64 = math.log1p %arg_f64 : f64
 277     // CHECK: llvm.call @__nv_log1p(%{{.*}}) : (f64) -> f64
 278     std.return %result32, %result64 : f32, f64
 279   }
 280 }
 281
 282 // -----
 283
 284 gpu.module @test_module {
 285   // CHECK: llvm.func @__nv_log2f(f32) -> f32
 286   // CHECK: llvm.func @__nv_log2(f64) -> f64
 287   // CHECK-LABEL: func @gpu_log2
 288   builtin.func @gpu_log2(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
 289     %result32 = math.log2 %arg_f32 : f32
 290     // CHECK: llvm.call @__nv_log2f(%{{.*}}) : (f32) -> f32
 291     %result64 = math.log2 %arg_f64 : f64
 292     // CHECK: llvm.call @__nv_log2(%{{.*}}) : (f64) -> f64
 293     std.return %result32, %result64 : f32, f64
 294   }
 295 }
 296
 297 // -----
 298
 299 gpu.module @test_module {
 300   // CHECK: llvm.func @__nv_sinf(f32) -> f32
 301   // CHECK: llvm.func @__nv_sin(f64) -> f64
 302   // CHECK-LABEL: func @gpu_sin
 303   builtin.func @gpu_sin(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
 304     %result32 = math.sin %arg_f32 : f32
 305     // CHECK: llvm.call @__nv_sinf(%{{.*}}) : (f32) -> f32
 306     %result64 = math.sin %arg_f64 : f64
 307     // CHECK: llvm.call @__nv_sin(%{{.*}}) : (f64) -> f64
 308     std.return %result32, %result64 : f32, f64
 309   }
 310 }
 311
 312 // -----
 313
 314 gpu.module @test_module {
 315   // CHECK: llvm.func @__nv_tanhf(f32) -> f32
 316   // CHECK: llvm.func @__nv_tanh(f64) -> f64
 317   // CHECK-LABEL: func @gpu_tanh
 318   builtin.func @gpu_tanh(%arg_f16 : f16, %arg_f32 : f32, %arg_f64 : f64) -> (f16, f32, f64) {
 319     %result16 = math.tanh %arg_f16 : f16
 320     // CHECK: llvm.fpext %{{.*}} : f16 to f32
 321     // CHECK-NEXT: llvm.call @__nv_tanhf(%{{.*}}) : (f32) -> f32
 322     // CHECK-NEXT: llvm.fptrunc %{{.*}} : f32 to f16
 323     %result32 = math.tanh %arg_f32 : f32
 324     // CHECK: llvm.call @__nv_tanhf(%{{.*}}) : (f32) -> f32
 325     %result64 = math.tanh %arg_f64 : f64
 326     // CHECK: llvm.call @__nv_tanh(%{{.*}}) : (f64) -> f64
 327     std.return %result16, %result32, %result64 : f16, f32, f64
 328   }
 329 }
 330
 331 // -----
 332
 333 gpu.module @test_module {
 334   // CHECK: llvm.func @__nv_rsqrtf(f32) -> f32
 335   // CHECK: llvm.func @__nv_rsqrt(f64) -> f64
 336   // CHECK-LABEL: func @gpu_rsqrt
 337   builtin.func @gpu_rsqrt(%arg_f16 : f16, %arg_f32 : f32, %arg_f64 : f64)
 338       -> (f16, f32, f64) {
 339     %result16 = math.rsqrt %arg_f16 : f16
 340     // CHECK: llvm.fpext %{{.*}} : f16 to f32
 341     // CHECK-NEXT: llvm.call @__nv_rsqrtf(%{{.*}}) : (f32) -> f32
 342     // CHECK-NEXT: llvm.fptrunc %{{.*}} : f32 to f16
 343     %result32 = math.rsqrt %arg_f32 : f32
 344     // CHECK: llvm.call @__nv_rsqrtf(%{{.*}}) : (f32) -> f32
 345     %result64 = math.rsqrt %arg_f64 : f64
 346     // CHECK: llvm.call @__nv_rsqrt(%{{.*}}) : (f64) -> f64
 347     std.return %result16, %result32, %result64 : f16, f32, f64
 348   }
 349 }
 350
 351 // -----
 352
 353 gpu.module @test_module {
 354   // CHECK: llvm.func @__nv_sqrtf(f32) -> f32
 355   // CHECK: llvm.func @__nv_sqrt(f64) -> f64
 356   // CHECK-LABEL: func @gpu_sqrt
 357   builtin.func @gpu_sqrt(%arg_f16 : f16, %arg_f32 : f32, %arg_f64 : f64)
 358       -> (f16, f32, f64) {
 359     %result16 = math.sqrt %arg_f16 : f16
 360     // CHECK: llvm.fpext %{{.*}} : f16 to f32
 361     // CHECK-NEXT: llvm.call @__nv_sqrtf(%{{.*}}) : (f32) -> f32
 362     // CHECK-NEXT: llvm.fptrunc %{{.*}} : f32 to f16
 363     %result32 = math.sqrt %arg_f32 : f32
 364     // CHECK: llvm.call @__nv_sqrtf(%{{.*}}) : (f32) -> f32
 365     %result64 = math.sqrt %arg_f64 : f64
 366     // CHECK: llvm.call @__nv_sqrt(%{{.*}}) : (f64) -> f64
 367     std.return %result16, %result32, %result64 : f16, f32, f64
 368   }
 369 }
 370
 371 // -----
 372
 373 gpu.module @test_module {
 374   // CHECK: llvm.func @__nv_atanf(f32) -> f32
 375   // CHECK: llvm.func @__nv_atan(f64) -> f64
 376   // CHECK-LABEL: func @gpu_atan
 377   builtin.func @gpu_atan(%arg_f16 : f16, %arg_f32 : f32, %arg_f64 : f64)
 378       -> (f16, f32, f64) {
 379     %result16 = math.atan %arg_f16 : f16
 380     // CHECK: llvm.fpext %{{.*}} : f16 to f32
 381     // CHECK-NEXT: llvm.call @__nv_atanf(%{{.*}}) : (f32) -> f32
 382     // CHECK-NEXT: llvm.fptrunc %{{.*}} : f32 to f16
 383     %result32 = math.atan %arg_f32 : f32
 384     // CHECK: llvm.call @__nv_atanf(%{{.*}}) : (f32) -> f32
 385     %result64 = math.atan %arg_f64 : f64
 386     // CHECK: llvm.call @__nv_atan(%{{.*}}) : (f64) -> f64
 387     std.return %result16, %result32, %result64 : f16, f32, f64
 388   }
 389 }
 390
 391 // -----
 392
 393 gpu.module @test_module {
 394   // CHECK: llvm.func @__nv_atan2f(f32, f32) -> f32
 395   // CHECK: llvm.func @__nv_atan2(f64, f64) -> f64
 396   // CHECK-LABEL: func @gpu_atan2
 397   builtin.func @gpu_atan2(%arg_f16 : f16, %arg_f32 : f32, %arg_f64 : f64)
 398       -> (f16, f32, f64) {
 399     %result16 = math.atan2 %arg_f16, %arg_f16 : f16
 400     // CHECK: llvm.fpext %{{.*}} : f16 to f32
 401     // CHECK: llvm.fpext %{{.*}} : f16 to f32
 402     // CHECK-NEXT: llvm.call @__nv_atan2f(%{{.*}}) : (f32, f32) -> f32
 403     // CHECK-NEXT: llvm.fptrunc %{{.*}} : f32 to f16
 404     %result32 = math.atan2 %arg_f32, %arg_f32 : f32
 405     // CHECK: llvm.call @__nv_atan2f(%{{.*}}) : (f32, f32) -> f32
 406     %result64 = math.atan2 %arg_f64, %arg_f64 : f64
 407     // CHECK: llvm.call @__nv_atan2(%{{.*}}) : (f64, f64) -> f64
 408     std.return %result16, %result32, %result64 : f16, f32, f64
 409   }
 410 }
 411
 412 // -----
 413
 414 // Test that we handled properly operation with SymbolTable other than module op
 415 gpu.module @test_module {
 416   "test.symbol_scope"() ({
 417   // CHECK: test.symbol_scope
 418   // CHECK: llvm.func @__nv_expf(f32) -> f32
 419   // CHECK: llvm.func @__nv_exp(f64) -> f64
 420   // CHECK-LABEL: func @gpu_exp
 421     builtin.func @gpu_exp(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
 422       %result32 = math.exp %arg_f32 : f32
 423       // CHECK: llvm.call @__nv_expf(%{{.*}}) : (f32) -> f32
 424       %result64 = math.exp %arg_f64 : f64
 425       // CHECK: llvm.call @__nv_exp(%{{.*}}) : (f64) -> f64
 426       std.return %result32, %result64 : f32, f64
 427     }
 428     "test.finish" () : () -> ()
 429   }) : () -> ()
 430 }
 431
 432 // -----
 433
 434 gpu.module @test_module {
 435   // CHECK: llvm.func @__nv_expm1f(f32) -> f32
 436   // CHECK: llvm.func @__nv_expm1(f64) -> f64
 437   // CHECK-LABEL: func @gpu_expm1
 438   builtin.func @gpu_expm1(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
 439     %result32 = math.expm1 %arg_f32 : f32
 440     // CHECK: llvm.call @__nv_expm1f(%{{.*}}) : (f32) -> f32
 441     %result64 = math.expm1 %arg_f64 : f64
 442     // CHECK: llvm.call @__nv_expm1(%{{.*}}) : (f64) -> f64
 443     std.return %result32, %result64 : f32, f64
 444   }
 445 }
 446
 447 // -----
 448
 449 gpu.module @test_module {
 450   // CHECK: llvm.func @__nv_powf(f32, f32) -> f32
 451   // CHECK: llvm.func @__nv_pow(f64, f64) -> f64
 452   // CHECK-LABEL: func @gpu_pow
 453   builtin.func @gpu_pow(%arg_f32 : f32, %arg_f64 : f64) -> (f32, f64) {
 454     %result32 = math.powf %arg_f32, %arg_f32 : f32
 455     // CHECK: llvm.call @__nv_powf(%{{.*}}, %{{.*}}) : (f32, f32) -> f32
 456     %result64 = math.powf %arg_f64, %arg_f64 : f64
 457     // CHECK: llvm.call @__nv_pow(%{{.*}}, %{{.*}}) : (f64, f64) -> f64
 458     std.return %result32, %result64 : f32, f64
 459   }
 460 }
 461
 462 // -----
 463
 464 gpu.module @test_module {
 465   // CHECK-LABEL: @kernel_func
 466   // CHECK: attributes
 467   // CHECK: gpu.kernel
 468   // CHECK: nvvm.kernel
 469   gpu.func @kernel_func() kernel {
 470     gpu.return
 471   }
 472 }