53 lines · plain
1// RUN: %clang_cc1 -std=c++11 -fsyntax-only -triple nvptx-unknown-unknown -target-cpu sm_75 -verify %s2 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 19int 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 25int nonconstint = 256;26__launch_bounds__(nonconstint) void TestNonConstInt(void); // expected-error {{'launch_bounds' attribute requires parameter 0 to be an integer constant}}27 28const int constint = 512;29__launch_bounds__(constint) void TestConstInt(void);30__launch_bounds__(constint * 2 + 3) void TestConstIntExpr(void);31 32template <int a, int b> __launch_bounds__(a, b) void TestTemplate2Args(void) {}33template void TestTemplate2Args<128,7>(void);34 35template <int a> __launch_bounds__(a) void TestTemplate1Arg(void) {}36template void TestTemplate1Arg<128>(void);37 38template <class a>39__launch_bounds__(a) void TestTemplate1ArgClass(void) {} // expected-error {{'a' does not refer to a value}}40// expected-note@-2 {{declared here}}41 42template <int a, int b, int c>43__launch_bounds__(a + b, c + constint) void TestTemplateExpr(void) {}44template void TestTemplateExpr<128+constint, 3, 7>(void);45 46template <int... Args>47__launch_bounds__(Args) void TestTemplateVariadicArgs(void) {} // expected-error {{expression contains unexpanded parameter pack 'Args'}}48 49template <int... Args>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