| 1 | // RUN: %clang_cc1 -fsyntax-only -Wundefined-internal -verify %s |
|---|
| 2 | // RUN: %clang_cc1 -fsyntax-only -Wundefined-internal -fcuda-is-device -verify %s |
|---|
| 3 | |
|---|
| 4 | // RUN: %clang_cc1 -fsyntax-only -Wundefined-internal -fgpu-rdc -verify=rdc %s |
|---|
| 5 | // RUN: %clang_cc1 -fsyntax-only -Wundefined-internal -fcuda-is-device -fgpu-rdc -verify=rdc %s |
|---|
| 6 | |
|---|
| 7 | // Most of these declarations are fine in separate compilation mode. |
|---|
| 8 | |
|---|
| 9 | #include "Inputs/cuda.h" |
|---|
| 10 | |
|---|
| 11 | __device__ void foo() { |
|---|
| 12 | extern __shared__ int x; // expected-error {{__shared__ variable 'x' cannot be 'extern'}} |
|---|
| 13 | extern __shared__ int arr[]; // ok |
|---|
| 14 | extern __shared__ int arr0[0]; // expected-error {{__shared__ variable 'arr0' cannot be 'extern'}} |
|---|
| 15 | extern __shared__ int arr1[1]; // expected-error {{__shared__ variable 'arr1' cannot be 'extern'}} |
|---|
| 16 | extern __shared__ int* ptr ; // expected-error {{__shared__ variable 'ptr' cannot be 'extern'}} |
|---|
| 17 | } |
|---|
| 18 | |
|---|
| 19 | __host__ __device__ void bar() { |
|---|
| 20 | extern __shared__ int arr[]; // ok |
|---|
| 21 | extern __shared__ int arr0[0]; // expected-error {{__shared__ variable 'arr0' cannot be 'extern'}} |
|---|
| 22 | extern __shared__ int arr1[1]; // expected-error {{__shared__ variable 'arr1' cannot be 'extern'}} |
|---|
| 23 | extern __shared__ int* ptr ; // expected-error {{__shared__ variable 'ptr' cannot be 'extern'}} |
|---|
| 24 | } |
|---|
| 25 | |
|---|
| 26 | extern __shared__ int global; // expected-error {{__shared__ variable 'global' cannot be 'extern'}} |
|---|
| 27 | extern __shared__ int global_arr[]; // ok |
|---|
| 28 | extern __shared__ int global_arr1[1]; // expected-error {{__shared__ variable 'global_arr1' cannot be 'extern'}} |
|---|
| 29 | |
|---|
| 30 | // Check that, iff we're not in rdc mode, extern __shared__ can appear in an |
|---|
| 31 | // anonymous namespace / in a static function without generating a warning |
|---|
| 32 | // about a variable with internal linkage but no definition |
|---|
| 33 | // (-Wundefined-internal). |
|---|
| 34 | namespace { |
|---|
| 35 | extern __shared__ int global_arr[]; // rdc-warning {{has internal linkage but is not defined}} |
|---|
| 36 | __global__ void in_anon_ns() { |
|---|
| 37 | extern __shared__ int local_arr[]; // rdc-warning {{has internal linkage but is not defined}} |
|---|
| 38 | |
|---|
| 39 | // Touch arrays to generate the warning. |
|---|
| 40 | local_arr[0] = 0; // rdc-note {{used here}} |
|---|
| 41 | global_arr[0] = 0; // rdc-note {{used here}} |
|---|
| 42 | } |
|---|
| 43 | } // namespace |
|---|
| 44 | |
|---|