1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
|
// REQUIRES: amdgpu-registered-target
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip %s \
// RUN: -emit-llvm -o - | FileCheck -check-prefix=DEV %s
// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -x hip %s \
// RUN: -emit-llvm -o - | FileCheck -check-prefix=HOST %s
// Negative tests.
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip %s \
// RUN: -emit-llvm -o - | FileCheck -check-prefix=DEV-NEG %s
#include "Inputs/cuda.h"
// Test const var initialized with address of a const var.
// Both are promoted to device side.
// DEV-DAG: @_ZN5Test1L1aE = internal addrspace(4) constant i32 1
// DEV-DAG: @_ZN5Test11B2p1E = addrspace(4) externally_initialized constant i32* addrspacecast (i32 addrspace(4)* @_ZN5Test1L1aE to i32*)
// DEV-DAG: @_ZN5Test11B2p2E = addrspace(4) externally_initialized constant i32* addrspacecast (i32 addrspace(4)* @_ZN5Test1L1aE to i32*)
// DEV-DAG: @_ZN5Test12b2E = addrspace(1) externally_initialized global i32 1
// HOST-DAG: @_ZN5Test1L1aE = internal constant i32 1
// HOST-DAG: @_ZN5Test11B2p1E = constant i32* @_ZN5Test1L1aE
// HOST-DAG: @_ZN5Test11B2p2E = internal constant i32* undef
// HOST-DAG: @_ZN5Test12b1E = global i32 1
// HOST-DAG: @_ZN5Test12b2E = internal global i32 undef
namespace Test1 {
const int a = 1;
struct B {
static const int *const p1;
static __device__ const int *const p2;
};
const int *const B::p1 = &a;
__device__ const int *const B::p2 = &a;
int b1 = B::p1 == B::p2;
__device__ int b2 = B::p1 == B::p2;
}
// Test const var initialized with address of a non-cost var.
// Neither is promoted to device side.
// DEV-NEG-NOT: @_ZN5Test2L1aE
// DEV-NEG-NOT: @_ZN5Test21B1pE
// HOST-DAG: @_ZN5Test21aE = global i32 1
// HOST-DAG: @_ZN5Test21B1pE = constant i32* @_ZN5Test21aE
namespace Test2 {
int a = 1;
struct B {
static int *const p;
};
int *const B::p = &a;
}
|
