clang/test/CodeGenCUDA/launch-bounds.cu

   1 // RUN: %clang_cc1 %s -triple nvptx-unknown-unknown -fcuda-is-device -emit-llvm -o - | FileCheck %s
   2
   3 #include "Inputs/cuda.h"
   4
   5 #define MAX_THREADS_PER_BLOCK 256
   6 #define MIN_BLOCKS_PER_MP     2
   7
   8 // Test both max threads per block and Min cta per sm.
   9 extern "C" {
  10 __global__ void
  11 __launch_bounds__( MAX_THREADS_PER_BLOCK, MIN_BLOCKS_PER_MP )
  12 Kernel1()
  13 {
  14 }
  15 }
  16
  17 // CHECK: !{{[0-9]+}} = !{ptr @Kernel1, !"maxntidx", i32 256}
  18 // CHECK: !{{[0-9]+}} = !{ptr @Kernel1, !"minctasm", i32 2}
  19
  20 // Test only max threads per block. Min cta per sm defaults to 0, and
  21 // CodeGen doesn't output a zero value for minctasm.
  22 extern "C" {
  23 __global__ void
  24 __launch_bounds__( MAX_THREADS_PER_BLOCK )
  25 Kernel2()
  26 {
  27 }
  28 }
  29
  30 // CHECK: !{{[0-9]+}} = !{ptr @Kernel2, !"maxntidx", i32 256}
  31
  32 template <int max_threads_per_block>
  33 __global__ void
  34 __launch_bounds__(max_threads_per_block)
  35 Kernel3()
  36 {
  37 }
  38
  39 template __global__ void Kernel3<MAX_THREADS_PER_BLOCK>();
  40 // CHECK: !{{[0-9]+}} = !{ptr @{{.*}}Kernel3{{.*}}, !"maxntidx", i32 256}
  41
  42 template <int max_threads_per_block, int min_blocks_per_mp>
  43 __global__ void
  44 __launch_bounds__(max_threads_per_block, min_blocks_per_mp)
  45 Kernel4()
  46 {
  47 }
  48 template __global__ void Kernel4<MAX_THREADS_PER_BLOCK, MIN_BLOCKS_PER_MP>();
  49
  50 // CHECK: !{{[0-9]+}} = !{ptr @{{.*}}Kernel4{{.*}}, !"maxntidx", i32 256}
  51 // CHECK: !{{[0-9]+}} = !{ptr @{{.*}}Kernel4{{.*}}, !"minctasm", i32 2}
  52
  53 const int constint = 100;
  54 template <int max_threads_per_block, int min_blocks_per_mp>
  55 __global__ void
  56 __launch_bounds__(max_threads_per_block + constint,
  57                   min_blocks_per_mp + max_threads_per_block)
  58 Kernel5()
  59 {
  60 }
  61 template __global__ void Kernel5<MAX_THREADS_PER_BLOCK, MIN_BLOCKS_PER_MP>();
  62
  63 // CHECK: !{{[0-9]+}} = !{ptr @{{.*}}Kernel5{{.*}}, !"maxntidx", i32 356}
  64 // CHECK: !{{[0-9]+}} = !{ptr @{{.*}}Kernel5{{.*}}, !"minctasm", i32 258}
  65
  66 // Make sure we don't emit negative launch bounds values.
  67 __global__ void
  68 __launch_bounds__( -MAX_THREADS_PER_BLOCK, MIN_BLOCKS_PER_MP )
  69 Kernel6()
  70 {
  71 }
  72 // CHECK-NOT: !{{[0-9]+}} = !{ptr @{{.*}}Kernel6{{.*}}, !"maxntidx",
  73 // CHECK:     !{{[0-9]+}} = !{ptr @{{.*}}Kernel6{{.*}}, !"minctasm",
  74
  75 __global__ void
  76 __launch_bounds__( MAX_THREADS_PER_BLOCK, -MIN_BLOCKS_PER_MP )
  77 Kernel7()
  78 {
  79 }
  80 // CHECK:     !{{[0-9]+}} = !{ptr @{{.*}}Kernel7{{.*}}, !"maxntidx",
  81 // CHECK-NOT: !{{[0-9]+}} = !{ptr @{{.*}}Kernel7{{.*}}, !"minctasm",
  82
  83 const char constchar = 12;
  84 __global__ void __launch_bounds__(constint, constchar) Kernel8() {}
  85 // CHECK:     !{{[0-9]+}} = !{ptr @{{.*}}Kernel8{{.*}}, !"maxntidx", i32 100
  86 // CHECK:     !{{[0-9]+}} = !{ptr @{{.*}}Kernel8{{.*}}, !"minctasm", i32 12