132 lines
5.9 KiB
C++
132 lines
5.9 KiB
C++
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --check-globals all --version 6
|
|
// REQUIRES: amdgpu-registered-target
|
|
// RUN: %clang_cc1 -x hip -triple x86_64-unknown-linux-gnu -aux-triple amdgcn-amd-amdhsa -emit-llvm-bc %s -o %t-host.bc
|
|
// RUN: %clang_cc1 -x hip -triple x86_64-unknown-linux-gnu -aux-triple amdgcn-amd-amdhsa -emit-llvm %s -o - | FileCheck %s --check-prefix=HOST
|
|
// 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=DEVICE
|
|
|
|
#define __device__ __attribute__((device))
|
|
#define __host__ __attribute__((host))
|
|
|
|
extern "C" {
|
|
//.
|
|
// HOST: @__HostVar = global i32 1, align 4
|
|
// HOST: @__hip_cuid_ = global i8 0
|
|
// HOST: @llvm.compiler.used = appending global [1 x ptr] [ptr @__hip_cuid_], section "llvm.metadata"
|
|
// HOST: @HostFunc = weak alias i32 (), ptr @__HostFunc
|
|
// HOST: @HostFunc_ = alias i32 (), ptr @__HostFunc
|
|
// HOST: @HostVar = weak alias i32, ptr @__HostVar
|
|
// HOST: @HostVar_ = alias i32, ptr @__HostVar
|
|
// HOST: @Two = weak alias i32 (), ptr @__Two
|
|
// HOST: @Two_ = alias i32 (), ptr @__Two
|
|
// HOST: @_Z5Threev = weak alias i32 (), ptr @_Z7__Threev
|
|
// HOST: @_Z6Three_v = alias i32 (), ptr @_Z7__Threev
|
|
// HOST: @_Z4Fourv = weak alias i32 (), ptr @_Z6__Fourv
|
|
// HOST: @_Z4Fourf = weak alias float (float), ptr @_Z6__Fourf
|
|
//.
|
|
// DEVICE: @__hip_cuid_ = addrspace(1) global i8 0
|
|
// DEVICE: @llvm.compiler.used = appending addrspace(1) global [1 x ptr] [ptr addrspacecast (ptr addrspace(1) @__hip_cuid_ to ptr)], section "llvm.metadata"
|
|
// DEVICE: @One = weak alias i32 (), ptr @__One
|
|
// DEVICE: @One_ = alias i32 (), ptr @__One
|
|
// DEVICE: @Two = weak alias i32 (), ptr @__Two
|
|
// DEVICE: @Two_ = alias i32 (), ptr @__Two
|
|
// DEVICE: @_Z5Threev = weak alias i32 (), ptr @_Z7__Threev
|
|
// DEVICE: @_Z6Three_v = alias i32 (), ptr @_Z7__Threev
|
|
// DEVICE: @_Z4Fourv = weak alias i32 (), ptr @_Z6__Fourv
|
|
// DEVICE: @_Z4Fourf = weak alias float (float), ptr @_Z6__Fourf
|
|
//.
|
|
// HOST-LABEL: define dso_local i32 @__HostFunc(
|
|
// HOST-SAME: ) #[[ATTR0:[0-9]+]] {
|
|
// HOST-NEXT: [[ENTRY:.*:]]
|
|
// HOST-NEXT: ret i32 42
|
|
//
|
|
int __HostFunc(void) { return 42; }
|
|
int __HostVar = 1;
|
|
int HostFunc(void) __attribute__((weak, alias("__HostFunc")));
|
|
int HostFunc_(void) __attribute__((alias("__HostFunc")));
|
|
extern int __attribute__((weak, alias("__HostVar"))) HostVar;
|
|
extern int __attribute__((alias("__HostVar"))) HostVar_;
|
|
|
|
// DEVICE-LABEL: define dso_local i32 @__One(
|
|
// DEVICE-SAME: ) #[[ATTR0:[0-9]+]] {
|
|
// DEVICE-NEXT: [[ENTRY:.*:]]
|
|
// DEVICE-NEXT: ret i32 1
|
|
//
|
|
__device__ int __One(void) { return 1; }
|
|
__device__ int One(void) __attribute__((weak, alias("__One")));
|
|
__device__ int One_(void) __attribute__((alias("__One")));
|
|
|
|
// HOST-LABEL: define dso_local i32 @__Two(
|
|
// HOST-SAME: ) #[[ATTR0]] {
|
|
// HOST-NEXT: [[ENTRY:.*:]]
|
|
// HOST-NEXT: ret i32 2
|
|
//
|
|
// DEVICE-LABEL: define dso_local i32 @__Two(
|
|
// DEVICE-SAME: ) #[[ATTR0]] {
|
|
// DEVICE-NEXT: [[ENTRY:.*:]]
|
|
// DEVICE-NEXT: ret i32 2
|
|
//
|
|
__host__ __device__ int __Two(void) { return 2; }
|
|
__host__ __device__ int Two(void) __attribute__((weak, alias("__Two")));
|
|
__host__ __device__ int Two_(void) __attribute__((alias("__Two")));
|
|
}
|
|
|
|
// HOST-LABEL: define linkonce_odr noundef i32 @_Z7__Threev(
|
|
// HOST-SAME: ) #[[ATTR0]] comdat {
|
|
// HOST-NEXT: [[ENTRY:.*:]]
|
|
// HOST-NEXT: ret i32 5
|
|
//
|
|
// DEVICE-LABEL: define linkonce_odr noundef i32 @_Z7__Threev(
|
|
// DEVICE-SAME: ) #[[ATTR0]] comdat {
|
|
// DEVICE-NEXT: [[ENTRY:.*:]]
|
|
// DEVICE-NEXT: ret i32 5
|
|
//
|
|
__host__ __device__ constexpr int __Three(void) { return 5; }
|
|
__host__ __device__ int Three(void) __attribute__((weak, alias("_Z7__Threev")));
|
|
__host__ __device__ int Three_(void) __attribute__((alias("_Z7__Threev")));
|
|
|
|
|
|
// HOST-LABEL: define dso_local noundef i32 @_Z6__Fourv(
|
|
// HOST-SAME: ) #[[ATTR0]] {
|
|
// HOST-NEXT: [[ENTRY:.*:]]
|
|
// HOST-NEXT: ret i32 2
|
|
//
|
|
// DEVICE-LABEL: define dso_local noundef i32 @_Z6__Fourv(
|
|
// DEVICE-SAME: ) #[[ATTR0]] {
|
|
// DEVICE-NEXT: [[ENTRY:.*:]]
|
|
// DEVICE-NEXT: ret i32 2
|
|
//
|
|
__host__ __device__ int __Four(void) { return 2; }
|
|
// HOST-LABEL: define dso_local noundef float @_Z6__Fourf(
|
|
// HOST-SAME: float noundef [[F:%.*]]) #[[ATTR0]] {
|
|
// HOST-NEXT: [[ENTRY:.*:]]
|
|
// HOST-NEXT: [[F_ADDR:%.*]] = alloca float, align 4
|
|
// HOST-NEXT: store float [[F]], ptr [[F_ADDR]], align 4
|
|
// HOST-NEXT: [[TMP0:%.*]] = load float, ptr [[F_ADDR]], align 4
|
|
// HOST-NEXT: [[MUL:%.*]] = fmul contract float 2.000000e+00, [[TMP0]]
|
|
// HOST-NEXT: ret float [[MUL]]
|
|
//
|
|
// DEVICE-LABEL: define dso_local noundef float @_Z6__Fourf(
|
|
// DEVICE-SAME: float noundef [[F:%.*]]) #[[ATTR0]] {
|
|
// DEVICE-NEXT: [[ENTRY:.*:]]
|
|
// DEVICE-NEXT: [[F_ADDR:%.*]] = alloca float, align 4, addrspace(5)
|
|
// DEVICE-NEXT: [[F_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[F_ADDR]] to ptr
|
|
// DEVICE-NEXT: store float [[F]], ptr [[F_ADDR_ASCAST]], align 4
|
|
// DEVICE-NEXT: [[TMP0:%.*]] = load float, ptr [[F_ADDR_ASCAST]], align 4
|
|
// DEVICE-NEXT: [[MUL:%.*]] = fmul contract float 2.000000e+00, [[TMP0]]
|
|
// DEVICE-NEXT: ret float [[MUL]]
|
|
//
|
|
__host__ __device__ float __Four(float f) { return 2.0f * f; }
|
|
__host__ __device__ int Four(void) __attribute__((weak, alias("_Z6__Fourv")));
|
|
__host__ __device__ float Four(float f) __attribute__((weak, alias("_Z6__Fourf")));
|
|
//.
|
|
// HOST: attributes #[[ATTR0]] = { mustprogress noinline nounwind optnone "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+cx8,+mmx,+sse,+sse2,+x87" "uniform-work-group-size" }
|
|
//.
|
|
// DEVICE: attributes #[[ATTR0]] = { convergent mustprogress noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size" }
|
|
//.
|
|
// HOST: [[META0:![0-9]+]] = !{!"{{.*}}clang version {{.*}}"}
|
|
//.
|
|
// DEVICE: [[META0:![0-9]+]] = !{i32 1, !"amdhsa_code_object_version", i32 600}
|
|
// DEVICE: [[META1:![0-9]+]] = !{i32 1, !"amdgpu_printf_kind", !"hostcall"}
|
|
// DEVICE: [[META2:![0-9]+]] = !{!"{{.*}}clang version {{.*}}"}
|
|
//.
|