149 lines · plain
1// RUN: %clang_cc1 %s -triple nvptx-unknown-unknown -fcuda-is-device -emit-llvm -o - | FileCheck %s2// RUN: %clang_cc1 %s -triple nvptx-unknown-unknown -target-cpu sm_90 -DUSE_MAX_BLOCKS -fcuda-is-device -emit-llvm -o - | FileCheck -check-prefix=CHECK_MAX_BLOCKS %s3 4#include "Inputs/cuda.h"5 6#define MAX_THREADS_PER_BLOCK 2567#define MIN_BLOCKS_PER_MP 28#ifdef USE_MAX_BLOCKS9#define MAX_BLOCKS_PER_MP 410#endif11 12// CHECK: @Kernel1() #[[ATTR0:[0-9]+]]13// CHECK: @Kernel2() #[[ATTR1:[0-9]+]]14// CHECK: @{{.*}}Kernel3{{.*}}() #[[ATTR1]]15// CHECK: @{{.*}}Kernel4{{.*}}() #[[ATTR0]]16// CHECK: @{{.*}}Kernel5{{.*}}() #[[ATTR2:[0-9]+]]17// CHECK: @{{.*}}Kernel6{{.*}}() #[[ATTR3:[0-9]+]]18// CHECK: @{{.*}}Kernel7{{.*}}() #[[ATTR1]]19// CHECK: @{{.*}}Kernel8{{.*}}() #[[ATTR4:[0-9]+]]20 21// CHECK-DAG: attributes #[[ATTR0]] = {{{.*}} "nvvm.maxntid"="256" "nvvm.minctasm"="2" {{.*}}}22// CHECK-DAG: attributes #[[ATTR1]] = {{{.*}} "nvvm.maxntid"="256" {{.*}}}23// CHECK-DAG: attributes #[[ATTR2]] = {{{.*}} "nvvm.maxntid"="356" "nvvm.minctasm"="258" {{.*}}}24// CHECK-DAG: attributes #[[ATTR3]] = {{{.*}} "nvvm.minctasm"="2" {{.*}}}25// CHECK-DAG: attributes #[[ATTR4]] = {{{.*}} "nvvm.maxntid"="100" "nvvm.minctasm"="12" {{.*}}}26 27// CHECK_MAX_BLOCKS: @Kernel1_sm_90() #[[ATTR0:[0-9]+]]28// CHECK_MAX_BLOCKS: @{{.*}}Kernel4_sm_90{{.*}} #[[ATTR0]]29// CHECK_MAX_BLOCKS: @{{.*}}Kernel5_sm_90{{.*}} #[[ATTR1:[0-9]+]]30// CHECK_MAX_BLOCKS: @{{.*}}Kernel7_sm_90{{.*}} #[[ATTR2:[0-9]+]]31// CHECK_MAX_BLOCKS: @{{.*}}Kernel8_sm_90{{.*}} #[[ATTR3:[0-9]+]]32 33// CHECK_MAX_BLOCKS-DAG: attributes #[[ATTR0]] = {{{.*}} "nvvm.maxclusterrank"="4" "nvvm.maxntid"="256" "nvvm.minctasm"="2" {{.*}}}34// CHECK_MAX_BLOCKS-DAG: attributes #[[ATTR1]] = {{{.*}} "nvvm.maxclusterrank"="260" "nvvm.maxntid"="356" "nvvm.minctasm"="258" {{.*}}}35// CHECK_MAX_BLOCKS-DAG: attributes #[[ATTR2]] = {{{.*}} "nvvm.maxntid"="256" {{.*}}}36// CHECK_MAX_BLOCKS-DAG: attributes #[[ATTR3]] = {{{.*}} "nvvm.maxclusterrank"="14" "nvvm.maxntid"="100" "nvvm.minctasm"="12" {{.*}}}37 38// Test both max threads per block and Min cta per sm.39extern "C" {40__global__ void41__launch_bounds__( MAX_THREADS_PER_BLOCK, MIN_BLOCKS_PER_MP )42Kernel1()43{44}45}46 47#ifdef USE_MAX_BLOCKS48// Test max threads per block and min/max cta per sm.49extern "C" {50__global__ void51__launch_bounds__( MAX_THREADS_PER_BLOCK, MIN_BLOCKS_PER_MP, MAX_BLOCKS_PER_MP )52Kernel1_sm_90()53{54}55}56#endif // USE_MAX_BLOCKS57 58// Test only max threads per block. Min cta per sm defaults to 0, and59// CodeGen doesn't output a zero value for minctasm.60extern "C" {61__global__ void62__launch_bounds__( MAX_THREADS_PER_BLOCK )63Kernel2()64{65}66}67 68template <int max_threads_per_block>69__global__ void70__launch_bounds__(max_threads_per_block)71Kernel3()72{73}74 75template __global__ void Kernel3<MAX_THREADS_PER_BLOCK>();76 77template <int max_threads_per_block, int min_blocks_per_mp>78__global__ void79__launch_bounds__(max_threads_per_block, min_blocks_per_mp)80Kernel4()81{82}83template __global__ void Kernel4<MAX_THREADS_PER_BLOCK, MIN_BLOCKS_PER_MP>();84 85 86#ifdef USE_MAX_BLOCKS87template <int max_threads_per_block, int min_blocks_per_mp, int max_blocks_per_mp>88__global__ void89__launch_bounds__(max_threads_per_block, min_blocks_per_mp, max_blocks_per_mp)90Kernel4_sm_90()91{92}93template __global__ void Kernel4_sm_90<MAX_THREADS_PER_BLOCK, MIN_BLOCKS_PER_MP, MAX_BLOCKS_PER_MP>();94 95#endif //USE_MAX_BLOCKS96 97const int constint = 100;98template <int max_threads_per_block, int min_blocks_per_mp>99__global__ void100__launch_bounds__(max_threads_per_block + constint,101 min_blocks_per_mp + max_threads_per_block)102Kernel5()103{104}105template __global__ void Kernel5<MAX_THREADS_PER_BLOCK, MIN_BLOCKS_PER_MP>();106 107#ifdef USE_MAX_BLOCKS108 109template <int max_threads_per_block, int min_blocks_per_mp, int max_blocks_per_mp>110__global__ void111__launch_bounds__(max_threads_per_block + constint,112 min_blocks_per_mp + max_threads_per_block,113 max_blocks_per_mp + max_threads_per_block)114Kernel5_sm_90()115{116}117template __global__ void Kernel5_sm_90<MAX_THREADS_PER_BLOCK, MIN_BLOCKS_PER_MP, MAX_BLOCKS_PER_MP>();118 119#endif //USE_MAX_BLOCKS120 121// Make sure we don't emit negative launch bounds values.122__global__ void123__launch_bounds__( -MAX_THREADS_PER_BLOCK, MIN_BLOCKS_PER_MP )124Kernel6()125{126}127 128__global__ void129__launch_bounds__( MAX_THREADS_PER_BLOCK, -MIN_BLOCKS_PER_MP )130Kernel7()131{132}133 134#ifdef USE_MAX_BLOCKS135__global__ void136__launch_bounds__( MAX_THREADS_PER_BLOCK, -MIN_BLOCKS_PER_MP, -MAX_BLOCKS_PER_MP )137Kernel7_sm_90()138{139}140#endif // USE_MAX_BLOCKS141 142const char constchar = 12;143__global__ void __launch_bounds__(constint, constchar) Kernel8() {}144 145#ifdef USE_MAX_BLOCKS146const char constchar_2 = 14;147__global__ void __launch_bounds__(constint, constchar, constchar_2) Kernel8_sm_90() {}148#endif // USE_MAX_BLOCKS149