126 lines · cpp
1// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --check-globals all --version 62// REQUIRES: amdgpu-registered-target3// RUN: %clang_cc1 -x hip -triple x86_64-unknown-linux-gnu -aux-triple amdgcn-amd-amdhsa -emit-llvm-bc %s -o %t-host.bc4// RUN: %clang_cc1 -x hip -triple x86_64-unknown-linux-gnu -aux-triple amdgcn-amd-amdhsa -emit-llvm %s -o - | FileCheck %s --check-prefix=HOST5// RUN: %clang_cc1 -x hip -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-linux-gnu -emit-llvm %s -fcuda-is-device -o - | FileCheck %s --check-prefix=DEVICE6 7#define __device__ __attribute__((device))8#define __host__ __attribute__((host))9 10extern "C" {11//.12// HOST: @HostFunc = weak alias i32 (), ptr @__HostFunc13// HOST: @HostFunc_ = alias i32 (), ptr @__HostFunc14// HOST: @HostVar = weak alias i32, ptr @__HostVar15// HOST: @HostVar_ = alias i32, ptr @__HostVar16// HOST: @Two = weak alias i32 (), ptr @__Two17// HOST: @Two_ = alias i32 (), ptr @__Two18// HOST: @_Z5Threev = weak alias i32 (), ptr @_Z7__Threev19// HOST: @_Z6Three_v = alias i32 (), ptr @_Z7__Threev20// HOST: @_Z4Fourv = weak alias i32 (), ptr @_Z6__Fourv21// HOST: @_Z4Fourf = weak alias float (float), ptr @_Z6__Fourf22//.23// DEVICE: @One = weak alias i32 (), ptr @__One24// DEVICE: @One_ = alias i32 (), ptr @__One25// DEVICE: @Two = weak alias i32 (), ptr @__Two26// DEVICE: @Two_ = alias i32 (), ptr @__Two27// DEVICE: @_Z5Threev = weak alias i32 (), ptr @_Z7__Threev28// DEVICE: @_Z6Three_v = alias i32 (), ptr @_Z7__Threev29// DEVICE: @_Z4Fourv = weak alias i32 (), ptr @_Z6__Fourv30// DEVICE: @_Z4Fourf = weak alias float (float), ptr @_Z6__Fourf31//.32// HOST-LABEL: define dso_local i32 @__HostFunc(33// HOST-SAME: ) #[[ATTR0:[0-9]+]] {34// HOST-NEXT: [[ENTRY:.*:]]35// HOST-NEXT: ret i32 4236//37int __HostFunc(void) { return 42; }38int __HostVar = 1;39int HostFunc(void) __attribute__((weak, alias("__HostFunc")));40int HostFunc_(void) __attribute__((alias("__HostFunc")));41extern int __attribute__((weak, alias("__HostVar"))) HostVar;42extern int __attribute__((alias("__HostVar"))) HostVar_;43 44// DEVICE-LABEL: define dso_local i32 @__One(45// DEVICE-SAME: ) #[[ATTR0:[0-9]+]] {46// DEVICE-NEXT: [[ENTRY:.*:]]47// DEVICE-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)48// DEVICE-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr49// DEVICE-NEXT: ret i32 150//51__device__ int __One(void) { return 1; }52__device__ int One(void) __attribute__((weak, alias("__One")));53__device__ int One_(void) __attribute__((alias("__One")));54 55// HOST-LABEL: define dso_local i32 @__Two(56// HOST-SAME: ) #[[ATTR0]] {57// HOST-NEXT: [[ENTRY:.*:]]58// HOST-NEXT: ret i32 259//60// DEVICE-LABEL: define dso_local i32 @__Two(61// DEVICE-SAME: ) #[[ATTR0]] {62// DEVICE-NEXT: [[ENTRY:.*:]]63// DEVICE-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)64// DEVICE-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr65// DEVICE-NEXT: ret i32 266//67__host__ __device__ int __Two(void) { return 2; }68__host__ __device__ int Two(void) __attribute__((weak, alias("__Two")));69__host__ __device__ int Two_(void) __attribute__((alias("__Two")));70}71 72// HOST-LABEL: define linkonce_odr noundef i32 @_Z7__Threev(73// HOST-SAME: ) #[[ATTR0]] comdat {74// HOST-NEXT: [[ENTRY:.*:]]75// HOST-NEXT: ret i32 576//77// DEVICE-LABEL: define linkonce_odr noundef i32 @_Z7__Threev(78// DEVICE-SAME: ) #[[ATTR0]] comdat {79// DEVICE-NEXT: [[ENTRY:.*:]]80// DEVICE-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)81// DEVICE-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr82// DEVICE-NEXT: ret i32 583//84__host__ __device__ constexpr int __Three(void) { return 5; }85__host__ __device__ int Three(void) __attribute__((weak, alias("_Z7__Threev")));86__host__ __device__ int Three_(void) __attribute__((alias("_Z7__Threev")));87 88 89// HOST-LABEL: define dso_local noundef i32 @_Z6__Fourv(90// HOST-SAME: ) #[[ATTR0]] {91// HOST-NEXT: [[ENTRY:.*:]]92// HOST-NEXT: ret i32 293//94// DEVICE-LABEL: define dso_local noundef i32 @_Z6__Fourv(95// DEVICE-SAME: ) #[[ATTR0]] {96// DEVICE-NEXT: [[ENTRY:.*:]]97// DEVICE-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)98// DEVICE-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr99// DEVICE-NEXT: ret i32 2100//101__host__ __device__ int __Four(void) { return 2; }102// HOST-LABEL: define dso_local noundef float @_Z6__Fourf(103// HOST-SAME: float noundef [[F:%.*]]) #[[ATTR0]] {104// HOST-NEXT: [[ENTRY:.*:]]105// HOST-NEXT: [[F_ADDR:%.*]] = alloca float, align 4106// HOST-NEXT: store float [[F]], ptr [[F_ADDR]], align 4107// HOST-NEXT: [[TMP0:%.*]] = load float, ptr [[F_ADDR]], align 4108// HOST-NEXT: [[MUL:%.*]] = fmul contract float 2.000000e+00, [[TMP0]]109// HOST-NEXT: ret float [[MUL]]110//111// DEVICE-LABEL: define dso_local noundef float @_Z6__Fourf(112// DEVICE-SAME: float noundef [[F:%.*]]) #[[ATTR0]] {113// DEVICE-NEXT: [[ENTRY:.*:]]114// DEVICE-NEXT: [[RETVAL:%.*]] = alloca float, align 4, addrspace(5)115// DEVICE-NEXT: [[F_ADDR:%.*]] = alloca float, align 4, addrspace(5)116// DEVICE-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr117// DEVICE-NEXT: [[F_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[F_ADDR]] to ptr118// DEVICE-NEXT: store float [[F]], ptr [[F_ADDR_ASCAST]], align 4119// DEVICE-NEXT: [[TMP0:%.*]] = load float, ptr [[F_ADDR_ASCAST]], align 4120// DEVICE-NEXT: [[MUL:%.*]] = fmul contract float 2.000000e+00, [[TMP0]]121// DEVICE-NEXT: ret float [[MUL]]122//123__host__ __device__ float __Four(float f) { return 2.0f * f; }124__host__ __device__ int Four(void) __attribute__((weak, alias("_Z6__Fourv")));125__host__ __device__ float Four(float f) __attribute__((weak, alias("_Z6__Fourf")));126