xref: /aosp_15_r20/external/clang/test/SemaCUDA/launch_bounds.cu (revision 67e74705e28f6214e480b399dd47ea732279e315)
1*67e74705SXin Li // RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify %s
2*67e74705SXin Li 
3*67e74705SXin Li #include "Inputs/cuda.h"
4*67e74705SXin Li 
5*67e74705SXin Li __launch_bounds__(128, 7) void Test2Args(void);
6*67e74705SXin Li __launch_bounds__(128) void Test1Arg(void);
7*67e74705SXin Li 
8*67e74705SXin Li __launch_bounds__(0xffffffff) void TestMaxArg(void);
9*67e74705SXin Li __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*67e74705SXin Li __launch_bounds__(0x10000000000000000) void TestWayTooBigArg(void); // expected-error {{integer literal is too large to be represented in any integer type}}
11*67e74705SXin Li 
12*67e74705SXin Li __launch_bounds__(-128, 7) void TestNegArg1(void); // expected-warning {{'launch_bounds' attribute parameter 0 is negative and will be ignored}}
13*67e74705SXin Li __launch_bounds__(128, -7) void TestNegArg2(void); // expected-warning {{'launch_bounds' attribute parameter 1 is negative and will be ignored}}
14*67e74705SXin Li 
15*67e74705SXin Li __launch_bounds__(1, 2, 3) void Test3Args(void); // expected-error {{'launch_bounds' attribute takes no more than 2 arguments}}
16*67e74705SXin Li __launch_bounds__() void TestNoArgs(void); // expected-error {{'launch_bounds' attribute takes at least 1 argument}}
17*67e74705SXin Li 
18*67e74705SXin Li int TestNoFunction __launch_bounds__(128, 7); // expected-warning {{'launch_bounds' attribute only applies to functions and methods}}
19*67e74705SXin Li 
20*67e74705SXin Li __launch_bounds__(true) void TestBool(void);
21*67e74705SXin Li __launch_bounds__(128.0) void TestFP(void); // expected-error {{'launch_bounds' attribute requires parameter 0 to be an integer constant}}
22*67e74705SXin Li __launch_bounds__((void*)0) void TestNullptr(void); // expected-error {{'launch_bounds' attribute requires parameter 0 to be an integer constant}}
23*67e74705SXin Li 
24*67e74705SXin Li int nonconstint = 256;
25*67e74705SXin Li __launch_bounds__(nonconstint) void TestNonConstInt(void); // expected-error {{'launch_bounds' attribute requires parameter 0 to be an integer constant}}
26*67e74705SXin Li 
27*67e74705SXin Li const int constint = 512;
28*67e74705SXin Li __launch_bounds__(constint) void TestConstInt(void);
29*67e74705SXin Li __launch_bounds__(constint * 2 + 3) void TestConstIntExpr(void);
30*67e74705SXin Li 
__launch_bounds__(a,b)31*67e74705SXin Li template <int a, int b> __launch_bounds__(a, b) void TestTemplate2Args(void) {}
32*67e74705SXin Li template void TestTemplate2Args<128,7>(void);
33*67e74705SXin Li 
__launch_bounds__(a)34*67e74705SXin Li template <int a> __launch_bounds__(a) void TestTemplate1Arg(void) {}
35*67e74705SXin Li template void TestTemplate1Arg<128>(void);
36*67e74705SXin Li 
37*67e74705SXin Li template <class a>
__launch_bounds__(a)38*67e74705SXin Li __launch_bounds__(a) void TestTemplate1ArgClass(void) {} // expected-error {{'a' does not refer to a value}}
39*67e74705SXin Li // expected-note@-2 {{declared here}}
40*67e74705SXin Li 
41*67e74705SXin Li template <int a, int b, int c>
TestTemplateExpr(void)42*67e74705SXin Li __launch_bounds__(a + b, c + constint) void TestTemplateExpr(void) {}
43*67e74705SXin Li template void TestTemplateExpr<128+constint, 3, 7>(void);
44*67e74705SXin Li 
45*67e74705SXin Li template <int... Args>
__launch_bounds__(Args)46*67e74705SXin Li __launch_bounds__(Args) void TestTemplateVariadicArgs(void) {} // expected-error {{expression contains unexpanded parameter pack 'Args'}}
47*67e74705SXin Li 
48*67e74705SXin Li template <int... Args>
TestTemplateVariadicArgs2(void)49*67e74705SXin Li __launch_bounds__(1, Args) void TestTemplateVariadicArgs2(void) {} // expected-error {{expression contains unexpanded parameter pack 'Args'}}
50