clang/test/CodeGenCUDA/record-layout.cu

   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=AST
   3 // RUN: cat %t | FileCheck --check-prefixes=CHECK,HOST %s
   4 // 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=AST
   7 // RUN: cat %t | FileCheck --check-prefixes=CHECK,DEV %s
   8
   9 #include "Inputs/cuda.h"
  10
  11 // AST: *** Dumping AST Record Layout
  12 // AST-LABEL:         0 | struct C
  13 // AST-NEXT:          0 |   struct A (base) (empty)
  14 // AST-NEXT:          1 |   struct B (base) (empty)
  15 // AST-NEXT:          4 |   int i
  16 // AST-NEXT:            | [sizeof=8, align=4,
  17 // AST-NEXT:            |  nvsize=8, nvalign=4]
  18
  19 // CHECK: %struct.C = type { [4 x i8], i32 }
  20
  21 struct A {};
  22 struct B {};
  23 struct C : A, B {
  24     int i;
  25 };
  26
  27 // AST: *** Dumping AST Record Layout
  28 // AST-LABEL:          0 | struct I
  29 // AST-NEXT:           0 |   (I vftable pointer)
  30 // AST-NEXT:           8 |   int i
  31 // AST-NEXT:             | [sizeof=16, align=8,
  32 // AST-NEXT:             |  nvsize=16, nvalign=8]
  33
  34 // AST: *** Dumping AST Record Layout
  35 // AST-LABEL:          0 | struct J
  36 // AST-NEXT:           0 |   struct I (primary base)
  37 // AST-NEXT:           0 |     (I vftable pointer)
  38 // AST-NEXT:           8 |     int i
  39 // AST-NEXT:          16 |   int j
  40 // 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 8
  52 // 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 8
  53 struct I {
  54     virtual void f() = 0;
  55     __device__ virtual void g() = 0;
  56     __device__ __host__ virtual void h() = 0;
  57     int i;
  58 };
  59
  60 struct 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 ptr
  70 // 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 1
  72 // DEV:  store i32 1, ptr %i, align 4
  73
  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 4
  81 // HOST:  %i = getelementptr inbounds nuw %struct.C, ptr %c, i32 0, i32 1
  82 // HOST:  store i32 11, ptr %i, align 4
  83
  84 void 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 ptr
  93 // DEV:  store ptr %j, ptr %j.addr.ascast, align 8
  94 // DEV:  %0 = load ptr, ptr %j.addr.ascast, align 8
  95 // DEV:  %i = getelementptr inbounds nuw %struct.I, ptr %0, i32 0, i32 1
  96 // DEV:  store i32 2, ptr %i, align 8
  97 // DEV:  %1 = load ptr, ptr %j.addr.ascast, align 8
  98 // DEV:  %j1 = getelementptr inbounds nuw %struct.J, ptr %1, i32 0, i32 1
  99 // DEV:  store i32 3, ptr %j1, align 8
 100 // DEV:  %2 = load ptr, ptr %j.addr.ascast, align 8
 101 // DEV:  %vtable = load ptr addrspace(1), ptr %2, align 8
 102 // DEV:  %vfn = getelementptr inbounds ptr addrspace(1), ptr addrspace(1) %vtable, i64 1
 103 // DEV:  %3 = load ptr addrspace(1), ptr addrspace(1) %vfn, align 8
 104 // DEV:  call addrspace(1) void %3(ptr noundef nonnull align 8 dereferenceable(24) %2)
 105 // DEV:  %4 = load ptr, ptr %j.addr.ascast, align 8
 106 // DEV:  %vtable2 = load ptr addrspace(1), ptr %4, align 8
 107 // DEV:  %vfn3 = getelementptr inbounds ptr addrspace(1), ptr addrspace(1) %vtable2, i64 2
 108 // DEV:  %5 = load ptr addrspace(1), ptr addrspace(1) %vfn3, align 8
 109 // 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 ptr
 121 // 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 8
 131 // HOST:  %i = getelementptr inbounds nuw %struct.I, ptr %0, i32 0, i32 1
 132 // HOST:  store i32 12, ptr %i, align 8
 133 // HOST:  %1 = load ptr, ptr %j.addr, align 8
 134 // HOST:  %j1 = getelementptr inbounds nuw %struct.J, ptr %1, i32 0, i32 1
 135 // HOST:  store i32 13, ptr %j1, align 8
 136 // HOST:  %2 = load ptr, ptr %j.addr, align 8
 137 // HOST:  %vtable = load ptr, ptr %2, align 8
 138 // HOST:  %vfn = getelementptr inbounds ptr, ptr %vtable, i64 0
 139 // HOST:  %3 = load ptr, ptr %vfn, align 8
 140 // HOST:  call void %3(ptr noundef nonnull align 8 dereferenceable(24) %2)
 141 // HOST:  %4 = load ptr, ptr %j.addr, align 8
 142 // HOST:  %vtable2 = load ptr, ptr %4, align 8
 143 // HOST:  %vfn3 = getelementptr inbounds ptr, ptr %vtable2, i64 2
 144 // HOST:  %5 = load ptr, ptr %vfn3, align 8
 145 // HOST:  call void %5(ptr noundef nonnull align 8 dereferenceable(24) %4)
 146
 147 void 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 8
 156 // 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
 159 void 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 8
 167 // HOST:  store ptr %this, ptr %this.addr, align 8
 168 // HOST:  %this1 = load ptr, ptr %this.addr, align 8
 169 // HOST:  %call = call noundef ptr @"??0I@@QEAA@XZ"(ptr noundef nonnull align 8 dereferenceable(16) %this1) #5
 170 // HOST:  store ptr @"??_7J@@6B@", ptr %this1, align 8
 171 // HOST:  ret ptr %this1
 172
 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 8
 175 // HOST:  store ptr %this, ptr %this.addr, align 8
 176 // HOST:  %this1 = load ptr, ptr %this.addr, align 8
 177 // HOST:  store ptr @"??_7I@@6B@", ptr %this1, align 8
 178 // HOST:  ret ptr %this1
 179
 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 ptr
 183 // DEV:  store ptr %this, ptr %this.addr.ascast, align 8
 184 // DEV:  %this1 = load ptr, ptr %this.addr.ascast, align 8
 185 // 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 ptr
 190 // DEV:  store ptr %this, ptr %this.addr.ascast, align 8
 191 // DEV:  %this1 = load ptr, ptr %this.addr.ascast, align 8
 192 // 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 8
 194
 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 ptr
 198 // DEV:  store ptr %this, ptr %this.addr.ascast, align 8
 199 // DEV:  %this1 = load ptr, ptr %this.addr.ascast, align 8
 200 // 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 8