201 lines · plain
1// RUN: %clang_cc1 -triple x86_64-pc-windows-msvc -fdump-record-layouts \2// RUN: -emit-llvm -o %t -xhip %s 2>&1 | FileCheck %s --check-prefix=AST3// RUN: cat %t | FileCheck --check-prefixes=CHECK,HOST %s4// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -target-cpu gfx1100 \5// RUN: -emit-llvm -fdump-record-layouts -aux-triple x86_64-pc-windows-msvc \6// RUN: -o %t -xhip %s | FileCheck %s --check-prefix=AST7// RUN: cat %t | FileCheck --check-prefixes=CHECK,DEV %s8 9#include "Inputs/cuda.h"10 11// AST: *** Dumping AST Record Layout12// AST-LABEL: 0 | struct C13// AST-NEXT: 0 | struct A (base) (empty)14// AST-NEXT: 1 | struct B (base) (empty)15// AST-NEXT: 4 | int i16// AST-NEXT: | [sizeof=8, align=4,17// AST-NEXT: | nvsize=8, nvalign=4]18 19// CHECK: %struct.C = type { [4 x i8], i32 }20 21struct A {};22struct B {};23struct C : A, B {24 int i;25};26 27// AST: *** Dumping AST Record Layout28// AST-LABEL: 0 | struct I29// AST-NEXT: 0 | (I vftable pointer)30// AST-NEXT: 8 | int i31// AST-NEXT: | [sizeof=16, align=8,32// AST-NEXT: | nvsize=16, nvalign=8]33 34// AST: *** Dumping AST Record Layout35// AST-LABEL: 0 | struct J36// AST-NEXT: 0 | struct I (primary base)37// AST-NEXT: 0 | (I vftable pointer)38// AST-NEXT: 8 | int i39// AST-NEXT: 16 | int j40// AST-NEXT: | [sizeof=24, align=8,41// AST-NEXT: | nvsize=24, nvalign=8]42 43// CHECK: %struct.I = type { ptr, i32 }44// CHECK: %struct.J = type { %struct.I, i32 }45 46// HOST: @0 = private unnamed_addr constant { [4 x ptr] } { [4 x ptr] [ptr @"??_R4J@@6B@", ptr @"?f@J@@UEAAXXZ", ptr null, ptr @"?h@J@@UEAAXXZ"] }, comdat($"??_7J@@6B@")47// HOST: @1 = private unnamed_addr constant { [4 x ptr] } { [4 x ptr] [ptr @"??_R4I@@6B@", ptr @_purecall, ptr null, ptr @_purecall] }, comdat($"??_7I@@6B@")48// HOST: @"??_7J@@6B@" = unnamed_addr alias ptr, getelementptr inbounds ({ [4 x ptr] }, ptr @0, i32 0, i32 0, i32 1)49// HOST: @"??_7I@@6B@" = unnamed_addr alias ptr, getelementptr inbounds ({ [4 x ptr] }, ptr @1, i32 0, i32 0, i32 1)50 51// DEV: @_ZTV1J = linkonce_odr unnamed_addr addrspace(1) constant { [5 x ptr addrspace(1)] } { [5 x ptr addrspace(1)] [ptr addrspace(1) null, ptr addrspace(1) null, ptr addrspace(1) null, ptr addrspace(1) addrspacecast (ptr @_ZN1J1gEv to ptr addrspace(1)), ptr addrspace(1) addrspacecast (ptr @_ZN1J1hEv to ptr addrspace(1))] }, comdat, align 852// DEV: @_ZTV1I = linkonce_odr unnamed_addr addrspace(1) constant { [5 x ptr addrspace(1)] } { [5 x ptr addrspace(1)] [ptr addrspace(1) null, ptr addrspace(1) null, ptr addrspace(1) null, ptr addrspace(1) addrspacecast (ptr @__cxa_pure_virtual to ptr addrspace(1)), ptr addrspace(1) addrspacecast (ptr @__cxa_pure_virtual to ptr addrspace(1))] }, comdat, align 853struct I {54 virtual void f() = 0;55 __device__ virtual void g() = 0;56 __device__ __host__ virtual void h() = 0;57 int i;58};59 60struct J : I {61 void f() override {}62 __device__ void g() override {}63 __device__ __host__ void h() override {}64 int j;65};66 67// DEV: define dso_local amdgpu_kernel void @_Z8C_kernel1C(ptr addrspace(4) noundef byref(%struct.C) align 4 %0)68// DEV: %coerce = alloca %struct.C, align 4, addrspace(5)69// DEV: %c = addrspacecast ptr addrspace(5) %coerce to ptr70// DEV: call void @llvm.memcpy.p0.p4.i64(ptr align 4 %c, ptr addrspace(4) align 4 %0, i64 8, i1 false)71// DEV: %i = getelementptr inbounds nuw %struct.C, ptr %c, i32 0, i32 172// DEV: store i32 1, ptr %i, align 473 74__global__ void C_kernel(C c)75{76 c.i = 1;77}78 79// HOST-LABEL: define dso_local void @"?test_C@@YAXXZ"()80// HOST: %c = alloca %struct.C, align 481// HOST: %i = getelementptr inbounds nuw %struct.C, ptr %c, i32 0, i32 182// HOST: store i32 11, ptr %i, align 483 84void test_C() {85 C c;86 c.i = 11;87 C_kernel<<<1, 1>>>(c);88}89 90// DEV: define dso_local void @_Z5J_devP1J(ptr noundef %j)91// DEV: %j.addr = alloca ptr, align 8, addrspace(5)92// DEV: %j.addr.ascast = addrspacecast ptr addrspace(5) %j.addr to ptr93// DEV: store ptr %j, ptr %j.addr.ascast, align 894// DEV: %0 = load ptr, ptr %j.addr.ascast, align 895// DEV: %i = getelementptr inbounds nuw %struct.I, ptr %0, i32 0, i32 196// DEV: store i32 2, ptr %i, align 897// DEV: %1 = load ptr, ptr %j.addr.ascast, align 898// DEV: %j1 = getelementptr inbounds nuw %struct.J, ptr %1, i32 0, i32 199// DEV: store i32 3, ptr %j1, align 8100// DEV: %2 = load ptr, ptr %j.addr.ascast, align 8101// DEV: %vtable = load ptr addrspace(1), ptr %2, align 8102// DEV: %vfn = getelementptr inbounds ptr addrspace(1), ptr addrspace(1) %vtable, i64 1103// DEV: %3 = load ptr addrspace(1), ptr addrspace(1) %vfn, align 8104// DEV: call addrspace(1) void %3(ptr noundef nonnull align 8 dereferenceable(24) %2)105// DEV: %4 = load ptr, ptr %j.addr.ascast, align 8106// DEV: %vtable2 = load ptr addrspace(1), ptr %4, align 8107// DEV: %vfn3 = getelementptr inbounds ptr addrspace(1), ptr addrspace(1) %vtable2, i64 2108// DEV: %5 = load ptr addrspace(1), ptr addrspace(1) %vfn3, align 8109// DEV: call addrspace(1) void %5(ptr noundef nonnull align 8 dereferenceable(24) %4)110 111__device__ void J_dev(J *j) {112 j->i = 2;113 j->j = 3;114 j->g();115 j->h();116}117 118// DEV: define dso_local amdgpu_kernel void @_Z8J_kernelv()119// DEV: %j = alloca %struct.J, align 8, addrspace(5)120// DEV: %j.ascast = addrspacecast ptr addrspace(5) %j to ptr121// DEV: call void @_ZN1JC1Ev(ptr noundef nonnull align 8 dereferenceable(24) %j.ascast)122// DEV: call void @_Z5J_devP1J(ptr noundef %j.ascast)123 124__global__ void J_kernel() {125 J j;126 J_dev(&j);127}128 129// HOST-LABEL: define dso_local void @"?J_host@@YAXPEAUJ@@@Z"(ptr noundef %j)130// HOST: %0 = load ptr, ptr %j.addr, align 8131// HOST: %i = getelementptr inbounds nuw %struct.I, ptr %0, i32 0, i32 1132// HOST: store i32 12, ptr %i, align 8133// HOST: %1 = load ptr, ptr %j.addr, align 8134// HOST: %j1 = getelementptr inbounds nuw %struct.J, ptr %1, i32 0, i32 1135// HOST: store i32 13, ptr %j1, align 8136// HOST: %2 = load ptr, ptr %j.addr, align 8137// HOST: %vtable = load ptr, ptr %2, align 8138// HOST: %vfn = getelementptr inbounds ptr, ptr %vtable, i64 0139// HOST: %3 = load ptr, ptr %vfn, align 8140// HOST: call void %3(ptr noundef nonnull align 8 dereferenceable(24) %2)141// HOST: %4 = load ptr, ptr %j.addr, align 8142// HOST: %vtable2 = load ptr, ptr %4, align 8143// HOST: %vfn3 = getelementptr inbounds ptr, ptr %vtable2, i64 2144// HOST: %5 = load ptr, ptr %vfn3, align 8145// HOST: call void %5(ptr noundef nonnull align 8 dereferenceable(24) %4)146 147void J_host(J *j) {148 j->i = 12;149 j->j = 13;150 j->f();151 j->h();152}153 154// HOST: define dso_local void @"?test_J@@YAXXZ"()155// HOST: %j = alloca %struct.J, align 8156// HOST: %call = call noundef ptr @"??0J@@QEAA@XZ"(ptr noundef nonnull align 8 dereferenceable(24) %j)157// HOST: call void @"?J_host@@YAXPEAUJ@@@Z"(ptr noundef %j)158 159void test_J() {160 J j;161 J_host(&j);162 J_kernel<<<1, 1>>>();163}164 165// HOST: define linkonce_odr dso_local noundef ptr @"??0J@@QEAA@XZ"(ptr noundef nonnull returned align 8 dereferenceable(24) %this)166// HOST: %this.addr = alloca ptr, align 8167// HOST: store ptr %this, ptr %this.addr, align 8168// HOST: %this1 = load ptr, ptr %this.addr, align 8169// HOST: %call = call noundef ptr @"??0I@@QEAA@XZ"(ptr noundef nonnull align 8 dereferenceable(16) %this1) #5170// HOST: store ptr @"??_7J@@6B@", ptr %this1, align 8171// HOST: ret ptr %this1172 173// HOST: define linkonce_odr dso_local noundef ptr @"??0I@@QEAA@XZ"(ptr noundef nonnull returned align 8 dereferenceable(16) %this)174// HOST: %this.addr = alloca ptr, align 8175// HOST: store ptr %this, ptr %this.addr, align 8176// HOST: %this1 = load ptr, ptr %this.addr, align 8177// HOST: store ptr @"??_7I@@6B@", ptr %this1, align 8178// HOST: ret ptr %this1179 180// DEV: define linkonce_odr void @_ZN1JC1Ev(ptr noundef nonnull align 8 dereferenceable(24) %this)181// DEV: %this.addr = alloca ptr, align 8, addrspace(5)182// DEV: %this.addr.ascast = addrspacecast ptr addrspace(5) %this.addr to ptr183// DEV: store ptr %this, ptr %this.addr.ascast, align 8184// DEV: %this1 = load ptr, ptr %this.addr.ascast, align 8185// DEV: call void @_ZN1JC2Ev(ptr noundef nonnull align 8 dereferenceable(24) %this1)186 187// DEV: define linkonce_odr void @_ZN1JC2Ev(ptr noundef nonnull align 8 dereferenceable(24) %this)188// DEV: %this.addr = alloca ptr, align 8, addrspace(5)189// DEV: %this.addr.ascast = addrspacecast ptr addrspace(5) %this.addr to ptr190// DEV: store ptr %this, ptr %this.addr.ascast, align 8191// DEV: %this1 = load ptr, ptr %this.addr.ascast, align 8192// DEV: call void @_ZN1IC2Ev(ptr noundef nonnull align 8 dereferenceable(16) %this1)193// DEV: store ptr addrspace(1) getelementptr inbounds inrange(-16, 24) ({ [5 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTV1J, i32 0, i32 0, i32 2), ptr %this1, align 8194 195// DEV: define linkonce_odr void @_ZN1IC2Ev(ptr noundef nonnull align 8 dereferenceable(16) %this)196// DEV: %this.addr = alloca ptr, align 8, addrspace(5)197// DEV: %this.addr.ascast = addrspacecast ptr addrspace(5) %this.addr to ptr198// DEV: store ptr %this, ptr %this.addr.ascast, align 8199// DEV: %this1 = load ptr, ptr %this.addr.ascast, align 8200// DEV: store ptr addrspace(1) getelementptr inbounds inrange(-16, 24) ({ [5 x ptr addrspace(1)] }, ptr addrspace(1) @_ZTV1I, i32 0, i32 0, i32 2), ptr %this1, align 8201