82 lines · plain
1// RUN: %clang_cc1 -triple amdgcn -fcuda-is-device \2// RUN: -emit-llvm -o - -x hip %s \3// RUN: | FileCheck -check-prefixes=DEV,NORDC %s4// RUN: %clang_cc1 -triple amdgcn -fcuda-is-device \5// RUN: -fgpu-rdc -cuid=abc -emit-llvm -o - -x hip %s \6// RUN: | FileCheck -check-prefixes=DEV,RDC %s7// RUN: %clang_cc1 -triple x86_64-unknown-gnu-linux \8// RUN: -emit-llvm -o - -x hip %s \9// RUN: | FileCheck -check-prefixes=HOST,NORDC-H %s10// RUN: %clang_cc1 -triple x86_64-unknown-gnu-linux \11// RUN: -fgpu-rdc -cuid=abc -emit-llvm -o - -x hip %s \12// RUN: | FileCheck -check-prefixes=HOST,RDC-H %s13// RUN: %clang_cc1 -triple nvptx -fcuda-is-device \14// RUN: -fgpu-rdc -cuid=abc -emit-llvm -o - %s \15// RUN: | FileCheck -check-prefixes=CUDA %s16 17#include "Inputs/cuda.h"18 19// DEV-DAG: @v1 = addrspace(1) externally_initialized global i32 020// NORDC-H-DAG: @v1 = internal global i32 undef21// RDC-H-DAG: @v1 = global i32 undef22__device__ int v1;23// DEV-DAG: @v2 = addrspace(4) externally_initialized constant i32 024// NORDC-H-DAG: @v2 = internal global i32 undef25// RDC-H-DAG: @v2 = global i32 undef26__constant__ int v2;27// DEV-DAG: @v3 = addrspace(1) externally_initialized global ptr addrspace(1) null28// NORDC-H-DAG: @v3 = internal externally_initialized global ptr null29// RDC-H-DAG: @v3 = externally_initialized global ptr null30#if __HIP__31__managed__ int v3;32#endif33 34// DEV-DAG: @ev1 = external addrspace(1) global i3235// HOST-DAG: @ev1 = external global i3236extern __device__ int ev1;37// DEV-DAG: @ev2 = external addrspace(4) global i3238// HOST-DAG: @ev2 = external global i3239extern __constant__ int ev2;40// DEV-DAG: @ev3 = external addrspace(1) externally_initialized global ptr addrspace(1)41// HOST-DAG: @ev3 = external externally_initialized global ptr42#if __HIP__43extern __managed__ int ev3;44#endif45 46// NORDC-DAG: @_ZL3sv1 = addrspace(1) externally_initialized global i32 047// RDC-DAG: @_ZL3sv1.static.[[HASH:.*]] = addrspace(1) externally_initialized global i32 048// HOST-DAG: @_ZL3sv1 = internal global i32 undef49// CUDA-DAG: @_ZL3sv1__static__[[HASH:.*]] = addrspace(1) externally_initialized global i32 050static __device__ int sv1;51// NORDC-DAG: @_ZL3sv2 = addrspace(4) externally_initialized constant i32 052// RDC-DAG: @_ZL3sv2.static.[[HASH]] = addrspace(4) externally_initialized constant i32 053// HOST-DAG: @_ZL3sv2 = internal global i32 undef54// CUDA-DAG: @_ZL3sv2__static__[[HASH]] = addrspace(4) externally_initialized constant i32 055static __constant__ int sv2;56// NORDC-DAG: @_ZL3sv3 = addrspace(1) externally_initialized global ptr addrspace(1) null57// RDC-DAG: @_ZL3sv3.static.[[HASH]] = addrspace(1) externally_initialized global ptr addrspace(1) null58// HOST-DAG: @_ZL3sv3 = internal externally_initialized global ptr null59#if __HIP__60static __managed__ int sv3;61#endif62 63__device__ __host__ int work(int *x);64 65__device__ __host__ int fun1() {66 return work(&ev1) + work(&ev2) + work(&sv1) + work(&sv2)67#if __HIP__68 + work(&ev3) + work(&sv3)69#endif70 ;71}72 73// HOST: hipRegisterVar({{.*}}@v174// HOST: hipRegisterVar({{.*}}@v275// HOST: hipRegisterManagedVar({{.*}}@v376// HOST-NOT: hipRegisterVar({{.*}}@ev177// HOST-NOT: hipRegisterVar({{.*}}@ev278// HOST-NOT: hipRegisterManagedVar({{.*}}@ev379// HOST: hipRegisterVar({{.*}}@_ZL3sv180// HOST: hipRegisterVar({{.*}}@_ZL3sv281// HOST: hipRegisterManagedVar({{.*}}@_ZL3sv382