xref: /llvm-project/clang/test/SemaCUDA/launch_bounds.cu (revision 3f8d4a8ef21dce88de0f79140556771d37944d19)
1 // RUN: %clang_cc1 -std=c++11 -fsyntax-only -triple nvptx-unknown-unknown -target-cpu sm_75 -verify %s
2 
3 #include "Inputs/cuda.h"
4 
5 __launch_bounds__(128, 7) void Test2Args(void);
6 __launch_bounds__(128) void Test1Arg(void);
7 
8 __launch_bounds__(0xffffffff) void TestMaxArg(void);
9 __launch_bounds__(0x100000000) void TestTooBigArg(void); // expected-error {{integer constant expression evaluates to value 4294967296 that cannot be represented in a 32-bit unsigned integer type}}
10 __launch_bounds__(0x10000000000000000) void TestWayTooBigArg(void); // expected-error {{integer literal is too large to be represented in any integer type}}
11 
12 __launch_bounds__(-128, 7) void TestNegArg1(void); // expected-warning {{'launch_bounds' attribute parameter 0 is negative and will be ignored}}
13 __launch_bounds__(128, -7) void TestNegArg2(void); // expected-warning {{'launch_bounds' attribute parameter 1 is negative and will be ignored}}
14 __launch_bounds__(128, 2, -8) void TestNegArg2(void); // expected-warning {{maxclusterrank requires sm_90 or higher, CUDA arch provided: sm_75, ignoring 'launch_bounds' attribute}}
15 
16 __launch_bounds__(1, 2, 3, 4) void Test4Args(void); // expected-error {{'launch_bounds' attribute takes no more than 3 arguments}}
17 __launch_bounds__() void TestNoArgs(void); // expected-error {{'launch_bounds' attribute takes at least 1 argument}}
18 
19 int TestNoFunction __launch_bounds__(128, 7); // expected-warning {{'launch_bounds' attribute only applies to Objective-C methods, functions, and function pointers}}
20 
21 __launch_bounds__(true) void TestBool(void);
22 __launch_bounds__(128.0) void TestFP(void); // expected-error {{'launch_bounds' attribute requires parameter 0 to be an integer constant}}
23 __launch_bounds__((void*)0) void TestNullptr(void); // expected-error {{'launch_bounds' attribute requires parameter 0 to be an integer constant}}
24 
25 int nonconstint = 256;
26 __launch_bounds__(nonconstint) void TestNonConstInt(void); // expected-error {{'launch_bounds' attribute requires parameter 0 to be an integer constant}}
27 
28 const int constint = 512;
29 __launch_bounds__(constint) void TestConstInt(void);
30 __launch_bounds__(constint * 2 + 3) void TestConstIntExpr(void);
31 
__launch_bounds__(a,b)32 template <int a, int b> __launch_bounds__(a, b) void TestTemplate2Args(void) {}
33 template void TestTemplate2Args<128,7>(void);
34 
__launch_bounds__(a)35 template <int a> __launch_bounds__(a) void TestTemplate1Arg(void) {}
36 template void TestTemplate1Arg<128>(void);
37 
38 template <class a>
__launch_bounds__(a)39 __launch_bounds__(a) void TestTemplate1ArgClass(void) {} // expected-error {{'a' does not refer to a value}}
40 // expected-note@-2 {{declared here}}
41 
42 template <int a, int b, int c>
TestTemplateExpr(void)43 __launch_bounds__(a + b, c + constint) void TestTemplateExpr(void) {}
44 template void TestTemplateExpr<128+constint, 3, 7>(void);
45 
46 template <int... Args>
__launch_bounds__(Args)47 __launch_bounds__(Args) void TestTemplateVariadicArgs(void) {} // expected-error {{expression contains unexpanded parameter pack 'Args'}}
48 
49 template <int... Args>
TestTemplateVariadicArgs2(void)50 __launch_bounds__(1, Args) void TestTemplateVariadicArgs2(void) {} // expected-error {{expression contains unexpanded parameter pack 'Args'}}
51 
52 __launch_bounds__(1, 2, 3) void Test3Args(void); // expected-warning {{maxclusterrank requires sm_90 or higher, CUDA arch provided: sm_75, ignoring 'launch_bounds' attribute}}
53