negative values and fix sm version to not rely on default

Author: jchlanda

clang/test/SemaCUDA/launch_bounds.cu

@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify %s
+// RUN: %clang_cc1 -std=c++11 -fsyntax-only -triple nvptx-unknown-unknown -target-cpu sm_75 -verify %s
 
 #include "Inputs/cuda.h"
 
@@ -11,6 +11,7 @@ __launch_bounds__(0x10000000000000000) void TestWayTooBigArg(void); // expected-
 
 __launch_bounds__(-128, 7) void TestNegArg1(void); // expected-warning {{'launch_bounds' attribute parameter 0 is negative and will be ignored}}
 __launch_bounds__(128, -7) void TestNegArg2(void); // expected-warning {{'launch_bounds' attribute parameter 1 is negative and will be ignored}}
+__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}}
 
 __launch_bounds__(1, 2, 3, 4) void Test4Args(void); // expected-error {{'launch_bounds' attribute takes no more than 3 arguments}}
 __launch_bounds__() void TestNoArgs(void); // expected-error {{'launch_bounds' attribute takes at least 1 argument}}
@@ -48,4 +49,4 @@ __launch_bounds__(Args) void TestTemplateVariadicArgs(void) {} // expected-error
 template <int... Args>
 __launch_bounds__(1, Args) void TestTemplateVariadicArgs2(void) {} // expected-error {{expression contains unexpanded parameter pack 'Args'}}
 
-__launch_bounds__(1, 2, 3) void Test3Args(void); // expected-warning {{maxclusterrank requires sm_90 or higher, CUDA arch provided: unknown, ignoring 'launch_bounds' attribute}}
+__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}}

clang/test/SemaCUDA/launch_bounds_sm_90.cu

@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -std=c++11 -fsyntax-only -triple nvptx-unknown-unknown -target-cpu sm_90  -verify %s
+// RUN: %clang_cc1 -std=c++11 -fsyntax-only -triple nvptx-unknown-unknown -target-cpu sm_90 -verify %s
 
 #include "Inputs/cuda.h"
 
@@ -12,7 +12,19 @@ __launch_bounds__(1, 1, 0x10000000000000000) void TestWayTooBigArg(void); // exp
 
 __launch_bounds__(-128, 7) void TestNegArg1(void); // expected-warning {{'launch_bounds' attribute parameter 0 is negative and will be ignored}}
 __launch_bounds__(128, -7) void TestNegArg2(void); // expected-warning {{'launch_bounds' attribute parameter 1 is negative and will be ignored}}
+__launch_bounds__(-128, 1, 7) void TestNegArg2(void); // expected-warning {{'launch_bounds' attribute parameter 0 is negative and will be ignored}}
+__launch_bounds__(128, -1, 7) void TestNegArg2(void); // expected-warning {{'launch_bounds' attribute parameter 1 is negative and will be ignored}}
 __launch_bounds__(128, 1, -7) void TestNegArg2(void); // expected-warning {{'launch_bounds' attribute parameter 2 is negative and will be ignored}}
+// expected-warning@20 {{'launch_bounds' attribute parameter 0 is negative and will be ignored}}
+// expected-warning@20 {{'launch_bounds' attribute parameter 1 is negative and will be ignored}}
+__launch_bounds__(-128, -1, 7) void TestNegArg2(void);
+// expected-warning@23 {{'launch_bounds' attribute parameter 0 is negative and will be ignored}}
+// expected-warning@23 {{'launch_bounds' attribute parameter 2 is negative and will be ignored}}
+__launch_bounds__(-128, 1, -7) void TestNegArg2(void);
+// expected-warning@27 {{'launch_bounds' attribute parameter 0 is negative and will be ignored}}
+// expected-warning@27 {{'launch_bounds' attribute parameter 1 is negative and will be ignored}}
+// expected-warning@27 {{'launch_bounds' attribute parameter 2 is negative and will be ignored}}
+__launch_bounds__(-128, -1, -7) void TestNegArg2(void);
 
 
 __launch_bounds__(1, 2, 3, 4) void Test4Args(void); // expected-error {{'launch_bounds' attribute takes no more than 3 arguments}}