This is an archive of the discontinued LLVM Phabricator instance.

Differential D143412

[HIP] Update test hip-header.hip
ClosedPublic

Authored by yaxunl on Feb 6 2023, 8:40 AM.

Details

Reviewers
tra
arsenm
Commits
rG4f01266954b1: [HIP] Update test hip-header.hip
Summary

remove -no-opaque-pointers

Diff Detail

Event Timeline

yaxunl created this revision.Feb 6 2023, 8:40 AM
Herald added a project: Restricted Project. · View Herald TranscriptFeb 6 2023, 8:40 AM
yaxunl requested review of this revision.Feb 6 2023, 8:40 AM
Herald added a subscriber: wdng. · View Herald TranscriptFeb 6 2023, 8:40 AM
arsenm accepted this revision.Feb 6 2023, 9:00 AM
This revision is now accepted and ready to land.Feb 6 2023, 9:00 AM
Harbormaster completed remote builds in B212116: Diff 495152.Feb 6 2023, 10:35 AM
Closed by commit rG4f01266954b1: [HIP] Update test hip-header.hip (authored by yaxunl). · Explain WhyFeb 6 2023, 1:33 PM
This revision was automatically updated to reflect the committed changes.
yaxunl added a commit: rG4f01266954b1: [HIP] Update test hip-header.hip.
Herald added a project: Restricted Project. · View Herald TranscriptFeb 6 2023, 1:33 PM

Revision Contents

PathSize
clang/
test/
Headers/
32 lines

Diff 495264

clang/test/Headers/hip-header.hip

// REQUIRES: amdgpu-registered-target // REQUIRES: amdgpu-registered-target
// RUN: %clang_cc1 -no-opaque-pointers -include __clang_hip_runtime_wrapper.h \ // RUN: %clang_cc1 -include __clang_hip_runtime_wrapper.h \
// RUN: -internal-isystem %S/../../lib/Headers/cuda_wrappers \ // RUN: -internal-isystem %S/../../lib/Headers/cuda_wrappers \
// RUN: -internal-isystem %S/Inputs/include \ // RUN: -internal-isystem %S/Inputs/include \
// RUN: -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown \ // RUN: -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown \
// RUN: -target-cpu gfx906 -emit-llvm %s -fcuda-is-device -o - \ // RUN: -target-cpu gfx906 -emit-llvm %s -fcuda-is-device -o - \
// RUN: -D__HIPCC_RTC__ | FileCheck -check-prefixes=CHECK,NOMALLOC %s // RUN: -D__HIPCC_RTC__ | FileCheck -check-prefixes=CHECK,NOMALLOC %s
// RUN: %clang_cc1 -no-opaque-pointers -include __clang_hip_runtime_wrapper.h \ // RUN: %clang_cc1 -include __clang_hip_runtime_wrapper.h \
// RUN: -internal-isystem %S/../../lib/Headers/cuda_wrappers \ // RUN: -internal-isystem %S/../../lib/Headers/cuda_wrappers \
// RUN: -internal-isystem %S/Inputs/include \ // RUN: -internal-isystem %S/Inputs/include \
// RUN: -include cmath \ // RUN: -include cmath \
// RUN: -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown \ // RUN: -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown \
// RUN: -target-cpu gfx906 -emit-llvm %s -fcuda-is-device -o - \ // RUN: -target-cpu gfx906 -emit-llvm %s -fcuda-is-device -o - \
// RUN: -D__HIPCC_RTC__ | FileCheck %s -check-prefixes=AMD_BOOL_RETURN // RUN: -D__HIPCC_RTC__ | FileCheck %s -check-prefixes=AMD_BOOL_RETURN
// RUN: %clang_cc1 -no-opaque-pointers -include __clang_hip_runtime_wrapper.h \ // RUN: %clang_cc1 -include __clang_hip_runtime_wrapper.h \
// RUN: -internal-isystem %S/../../lib/Headers/cuda_wrappers \ // RUN: -internal-isystem %S/../../lib/Headers/cuda_wrappers \
// RUN: -internal-isystem %S/Inputs/include \ // RUN: -internal-isystem %S/Inputs/include \
// RUN: -include cmath \ // RUN: -include cmath \
// RUN: -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown \ // RUN: -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown \
// RUN: -target-cpu gfx906 -emit-llvm %s -fcuda-is-device -o - \ // RUN: -target-cpu gfx906 -emit-llvm %s -fcuda-is-device -o - \
// RUN: -D__HIPCC_RTC__ -DUSE_ISNAN_WITH_INT_RETURN | FileCheck %s -check-prefixes=AMD_INT_RETURN // RUN: -D__HIPCC_RTC__ -DUSE_ISNAN_WITH_INT_RETURN | FileCheck %s -check-prefixes=AMD_INT_RETURN
// RUN: %clang_cc1 -no-opaque-pointers -include __clang_hip_runtime_wrapper.h \ // RUN: %clang_cc1 -include __clang_hip_runtime_wrapper.h \
// RUN: -internal-isystem %S/../../lib/Headers/cuda_wrappers \ // RUN: -internal-isystem %S/../../lib/Headers/cuda_wrappers \
// RUN: -internal-isystem %S/Inputs/include \ // RUN: -internal-isystem %S/Inputs/include \
// RUN: -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown \ // RUN: -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown \
// RUN: -target-cpu gfx906 -emit-llvm %s -fcuda-is-device -o - \ // RUN: -target-cpu gfx906 -emit-llvm %s -fcuda-is-device -o - \
// RUN: -D__HIPCC_RTC__ -std=c++14 | FileCheck -check-prefixes=CHECK,CXX14 %s // RUN: -D__HIPCC_RTC__ -std=c++14 | FileCheck -check-prefixes=CHECK,CXX14 %s
// RUN: %clang_cc1 -no-opaque-pointers -include __clang_hip_runtime_wrapper.h \ // RUN: %clang_cc1 -include __clang_hip_runtime_wrapper.h \
// RUN: -internal-isystem %S/../../lib/Headers/cuda_wrappers \ // RUN: -internal-isystem %S/../../lib/Headers/cuda_wrappers \
// RUN: -internal-isystem %S/Inputs/include \ // RUN: -internal-isystem %S/Inputs/include \
// RUN: -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown \ // RUN: -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown \
// RUN: -target-cpu gfx906 -emit-llvm %s -fcuda-is-device -o - \ // RUN: -target-cpu gfx906 -emit-llvm %s -fcuda-is-device -o - \
// RUN: -DHIP_VERSION_MAJOR=4 -DHIP_VERSION_MINOR=5 \ // RUN: -DHIP_VERSION_MAJOR=4 -DHIP_VERSION_MINOR=5 \
// RUN: -D__HIPCC_RTC__ | FileCheck -check-prefixes=CHECK,MALLOC %s // RUN: -D__HIPCC_RTC__ | FileCheck -check-prefixes=CHECK,MALLOC %s
// RUN: %clang_cc1 -no-opaque-pointers -include __clang_hip_runtime_wrapper.h \ // RUN: %clang_cc1 -include __clang_hip_runtime_wrapper.h \
// RUN: -internal-isystem %S/../../lib/Headers/cuda_wrappers \ // RUN: -internal-isystem %S/../../lib/Headers/cuda_wrappers \
// RUN: -internal-isystem %S/Inputs/include \ // RUN: -internal-isystem %S/Inputs/include \
// RUN: -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown \ // RUN: -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown \
// RUN: -target-cpu gfx906 -emit-llvm %s -fcuda-is-device -o - \ // RUN: -target-cpu gfx906 -emit-llvm %s -fcuda-is-device -o - \
// RUN: -DHIP_VERSION_MAJOR=4 -DHIP_VERSION_MINOR=5 \ // RUN: -DHIP_VERSION_MAJOR=4 -DHIP_VERSION_MINOR=5 \
// RUN: -fsanitize=address -disable-llvm-passes -D__HIPCC_RTC__ \ // RUN: -fsanitize=address -disable-llvm-passes -D__HIPCC_RTC__ \
// RUN: | FileCheck -check-prefixes=MALLOC-ASAN %s // RUN: | FileCheck -check-prefixes=MALLOC-ASAN %s
Show All 11 Lines
struct derived:base { struct derived:base {
__host__ __host__
__device__ __device__
virtual void pv() override {}; virtual void pv() override {};
}; };
__device__ void test_vf() { __device__ void test_vf() {
derived d; derived d;
} }
// CHECK: @_ZTV7derived = linkonce_odr unnamed_addr addrspace(1) constant { [4 x i8*] } { [4 x i8*] [i8* null, i8* null, i8* bitcast (void (%struct.derived*)* @_ZN7derived2pvEv to i8*), i8* bitcast (void ()* @__cxa_deleted_virtual to i8*)] }, comdat, align 8 // CHECK: @_ZTV7derived = linkonce_odr unnamed_addr addrspace(1) constant { [4 x ptr] } { [4 x ptr] [ptr null, ptr null, ptr @_ZN7derived2pvEv, ptr @__cxa_deleted_virtual] }, comdat, align 8
// CHECK: @_ZTV4base = linkonce_odr unnamed_addr addrspace(1) constant { [4 x i8*] } { [4 x i8*] [i8* null, i8* null, i8* bitcast (void ()* @__cxa_pure_virtual to i8*), i8* bitcast (void ()* @__cxa_deleted_virtual to i8*)] }, comdat, align 8 // CHECK: @_ZTV4base = linkonce_odr unnamed_addr addrspace(1) constant { [4 x ptr] } { [4 x ptr] [ptr null, ptr null, ptr @__cxa_pure_virtual, ptr @__cxa_deleted_virtual] }, comdat, align 8
// CHECK: define{{.*}}void @__cxa_pure_virtual() // CHECK: define{{.*}}void @__cxa_pure_virtual()
// CHECK: define{{.*}}void @__cxa_deleted_virtual() // CHECK: define{{.*}}void @__cxa_deleted_virtual()
struct Number { struct Number {
__device__ Number(float _x) : x(_x) {} __device__ Number(float _x) : x(_x) {}
float x; float x;
}; };
▲ Show 20 Lines • Show All 56 Lines • ▼ Show 20 Lines__device__ double test_isnan() {
r += isnan(d); r += isnan(d);
return r ; return r ;
} }
// Check that device malloc and free do not conflict with std headers. // Check that device malloc and free do not conflict with std headers.
#include <cstdlib> #include <cstdlib>
// CHECK-LABEL: define{{.*}}@_Z11test_malloc // CHECK-LABEL: define{{.*}}@_Z11test_malloc
// CHECK: call {{.*}}i8* @malloc(i64 // CHECK: call {{.*}}ptr @malloc(i64
// CHECK-LABEL: define weak {{.*}}i8* @malloc(i64 // CHECK-LABEL: define weak {{.*}}ptr @malloc(i64
// MALLOC: call i64 @__ockl_dm_alloc // MALLOC: call i64 @__ockl_dm_alloc
// NOMALLOC: call void @llvm.trap // NOMALLOC: call void @llvm.trap
// MALLOC-ASAN-LABEL: define weak {{.*}}i8* @malloc(i64 // MALLOC-ASAN-LABEL: define weak {{.*}}ptr @malloc(i64
// MALLOC-ASAN: call i8* @llvm.returnaddress(i32 0) // MALLOC-ASAN: call ptr @llvm.returnaddress(i32 0)
// MALLOC-ASAN: call i64 @__asan_malloc_impl(i64 {{.*}}, i64 {{.*}}) // MALLOC-ASAN: call i64 @__asan_malloc_impl(i64 {{.*}}, i64 {{.*}})
__device__ void test_malloc(void *a) { __device__ void test_malloc(void *a) {
a = malloc(42); a = malloc(42);
} }
// CHECK-LABEL: define{{.*}}@_Z9test_free // CHECK-LABEL: define{{.*}}@_Z9test_free
// CHECK: call {{.*}}void @free(i8* // CHECK: call {{.*}}void @free(ptr
// CHECK-LABEL: define weak {{.*}}void @free(i8* // CHECK-LABEL: define weak {{.*}}void @free(ptr
// MALLOC: call void @__ockl_dm_dealloc // MALLOC: call void @__ockl_dm_dealloc
// NOMALLOC: call void @llvm.trap // NOMALLOC: call void @llvm.trap
// MALLOC-ASAN-LABEL: define weak {{.*}}void @free(i8* // MALLOC-ASAN-LABEL: define weak {{.*}}void @free(ptr
// MALLOC-ASAN: call i8* @llvm.returnaddress(i32 0) // MALLOC-ASAN: call ptr @llvm.returnaddress(i32 0)
// MALLOC-ASAN: call void @__asan_free_impl(i64 {{.*}}, i64 {{.*}}) // MALLOC-ASAN: call void @__asan_free_impl(i64 {{.*}}, i64 {{.*}})
__device__ void test_free(void *a) { __device__ void test_free(void *a) {
free(a); free(a);
} }