107 lines · plain
1// RUN: %clang_cc1 -std=c++14 -triple amdgcn-amd-amdhsa \2// RUN: -fcuda-is-device -emit-llvm -o - -x hip %s | FileCheck %s3// RUN: %clang_cc1 -std=c++14 -triple amdgcn-amd-amdhsa \4// RUN: -fcuda-is-device -emit-llvm -o - -x hip %s | FileCheck -check-prefix=NEG %s5 6#include "Inputs/cuda.h"7 8struct A {9 int x;10};11 12// Check the situation of B<T> has empty ctor but B<int> has non-empty ctor.13// Make sure const B<int> variables are not promoted to constant variables.14template<typename T>15struct B {16 T x;17 B() {}18 B(T _x) { x = _x; }19 static const B<T> y;20};21 22template<>23struct B<int> {24 int x;25 B() { x = 1; }26 static const B<int> y;27};28 29template<typename T>30const B<T> B<T>::y;31 32const B<int> B<int>::y;33 34template<typename T>35T temp_fun(T x) {36 return B<T>::y.x;37}38 39// Check template variable with empty default ctor but non-empty initializer40// ctor is not promoted.41template<typename T>42const B<T> b = B<T>(-1);43 44constexpr int constexpr_var = 1;45constexpr A constexpr_struct{2};46constexpr A constexpr_array[4] = {0, 0, 0, 3};47constexpr char constexpr_str[] = "abcd";48const int const_var = 4;49const A const_struct{5};50const A const_array[] = {0, 0, 0, 6};51const char const_str[] = "xyz";52 53// Check const variables used by host only are not emitted.54const int var_host_only = 7;55 56// CHECK-DAG: @_ZL13constexpr_str.const = private unnamed_addr addrspace(4) constant [5 x i8] c"abcd\00"57// CHECK-DAG: @_ZL13constexpr_var = internal addrspace(4) constant i32 158// CHECK-DAG: @_ZL16constexpr_struct = internal addrspace(4) constant %struct.A { i32 2 }59// CHECK-DAG: @_ZL15constexpr_array = internal addrspace(4) constant [4 x %struct.A] [%struct.A zeroinitializer, %struct.A zeroinitializer, %struct.A zeroinitializer, %struct.A { i32 3 }]60// CHECK-DAG: @_ZL9const_var = internal addrspace(4) constant i32 461// CHECK-DAG: @_ZL12const_struct = internal addrspace(4) constant %struct.A { i32 5 }62// CHECK-DAG: @_ZL11const_array = internal addrspace(4) constant [4 x %struct.A] [%struct.A zeroinitializer, %struct.A zeroinitializer, %struct.A zeroinitializer, %struct.A { i32 6 }]63// CHECK-DAG: @_ZL9const_str = internal addrspace(4) constant [4 x i8] c"xyz\00"64 65// NEG-NOT: @_ZN1BIiE1yE66// NEG-NOT: @_Z1bIdE67// NEG-NOT: @_ZL13var_host_only68// NEG-NOT: {{^}}@{{.*}} = external69 70// CHECK-LABEL: define{{.*}}@_Z7dev_funPiPPKi71// CHECK: store i32 172// CHECK: store i32 273// CHECK: store i32 374// CHECK: load i8, ptr getelementptr {{.*}} @_ZL13constexpr_str.const75// CHECK: store i32 476// CHECK: store i32 577// CHECK: store i32 678// CHECK: load i8, ptr getelementptr {{.*}} @_ZL9const_str79// CHECK: store ptr {{.*}}@_ZL13constexpr_var80// CHECK: store ptr {{.*}} @_ZL16constexpr_struct81// CHECK: store ptr getelementptr {{.*}} @_ZL15constexpr_array82// CHECK: store ptr {{.*}}@_ZL9const_var83// CHECK: store ptr {{.*}} @_ZL12const_struct84// CHECK: store ptr getelementptr {{.*}} @_ZL11const_array85__device__ void dev_fun(int *out, const int **out2) {86 *out = constexpr_var;87 *out = constexpr_struct.x;88 *out = constexpr_array[3].x;89 *out = constexpr_str[3];90 *out = const_var;91 *out = const_struct.x;92 *out = const_array[3].x;93 *out = const_str[3];94 *out2 = &constexpr_var;95 *out2 = &constexpr_struct.x;96 *out2 = &constexpr_array[3].x;97 *out2 = &const_var;98 *out2 = &const_struct.x;99 *out2 = &const_array[3].x;100}101 102void fun() {103 temp_fun(1);104 (void) b<double>;105 (void) var_host_only;106}107