54 lines · plain
1// REQUIRES: amdgpu-registered-target2// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip %s \3// RUN: -std=c++11 -O3 -mllvm -amdgpu-internalize-symbols -emit-llvm -o - \4// RUN: -target-cpu gfx906 | FileCheck %s5// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip %s \6// RUN: -std=c++11 -O3 -mllvm -amdgpu-internalize-symbols -emit-llvm -o - \7// RUN: -target-cpu gfx906 | FileCheck -check-prefix=NEGCHK %s8 9#include "Inputs/cuda.h"10 11// AMDGPU internalize unused global variables for whole-program compilation12// (-fno-gpu-rdc for each TU, or -fgpu-rdc for LTO), which are then13// eliminated by global DCE. If there are invisible unused address space casts14// for global variables, these dead users need to be eliminated by global15// DCE before internalization. This test makes sure unused global variables16// are eliminated.17 18// CHECK-DAG: @v119__device__ int v1;20 21// CHECK-DAG: @v222__constant__ int v2;23 24// Check unused device/constant variables are eliminated.25 26// NEGCHK-NOT: @_ZL2v327constexpr int v3 = 1;28 29// Check managed variables are always kept.30 31// CHECK-DAG: @v432__managed__ int v4;33 34// Check used device/constant variables are not eliminated.35// CHECK-DAG: @u136__device__ int u1;37 38// CHECK-DAG: @u239__constant__ int u2;40 41// Check u3 is kept because its address is taken.42// CHECK-DAG: @_ZL2u343constexpr int u3 = 2;44 45// Check u4 is not kept because it is not ODR-use.46// NEGCHK-NOT: @_ZL2u447constexpr int u4 = 3;48 49__device__ int fun1(const int& x);50 51__global__ void kern1(int *x) {52 *x = u1 + u2 + fun1(u3) + u4;53}54