// REQUIRES: amdgpu-registered-target
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip %s \
// RUN:   -std=c++11 -O3 -mllvm -amdgpu-internalize-symbols -emit-llvm -o - \
// RUN:   | FileCheck %s

#include "Inputs/cuda.h"

// Check device variables used by neither host nor device functioins are not kept.

// CHECK-NOT: @v1
__device__ int v1;

// CHECK-NOT: @v2
__constant__ int v2;

// CHECK-NOT: @_ZL2v3
static __device__ int v3;

// Check device variables used by host functions are kept.

// CHECK-DAG: @u1
__device__ int u1;

// CHECK-DAG: @u2
__constant__ int u2;

// Check host-used static device var is in llvm.compiler.used.
// CHECK-DAG: @_ZL2u3
static __device__ int u3;

// Check device-used static device var is emitted but is not in llvm.compiler.used.
// CHECK-DAG: @_ZL2u4
static __device__ int u4;

// Check device variables with used attribute are always kept.
// CHECK-DAG: @u5
__device__ __attribute__((used)) int u5;

int fun1() {
  return u1 + u2 + u3;
}

__global__ void kern1(int **x) {
  *x = &u4;
}
// Check the exact list of variables to ensure @_ZL2u4 is not among them.
// CHECK: @llvm.compiler.used = {{[^@]*}} @_ZL2u3 {{[^@]*}} @u1 {{[^@]*}} @u2 {{[^@]*}} @u5