170 lines · plain
1// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -std=c++11 \2// RUN: -emit-llvm -o - -x hip %s | FileCheck \3// RUN: -check-prefixes=COMMON,DEV,NORDC-D %s4 5// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -std=c++11 \6// RUN: -emit-llvm -fgpu-rdc -cuid=abc -o - -x hip %s > %t.dev7// RUN: cat %t.dev | FileCheck -check-prefixes=COMMON,DEV,RDC-D %s8 9// RUN: %clang_cc1 -triple x86_64-gnu-linux -std=c++11 \10// RUN: -emit-llvm -o - -x hip %s | FileCheck \11// RUN: -check-prefixes=COMMON,HOST,NORDC %s12 13// RUN: %clang_cc1 -triple x86_64-gnu-linux -std=c++11 \14// RUN: -emit-llvm -fgpu-rdc -cuid=abc -o - -x hip %s > %t.host15// RUN: cat %t.host | FileCheck -check-prefixes=COMMON,HOST,RDC %s16 17// Check device and host compilation use the same postfix for static18// variable name.19 20// RUN: cat %t.dev %t.host | FileCheck -check-prefix=POSTFIX %s21 22#include "Inputs/cuda.h"23 24struct vec {25 float x,y,z;26};27 28// DEV-DAG: @x.managed = addrspace(1) externally_initialized global i32 1, align 429// DEV-DAG: @x = addrspace(1) externally_initialized global ptr addrspace(1) null30// NORDC-DAG: @x.managed = internal global i32 131// RDC-DAG: @x.managed = global i32 132// NORDC-DAG: @x = internal externally_initialized global ptr null33// RDC-DAG: @x = externally_initialized global ptr null34// HOST-DAG: @[[DEVNAMEX:[0-9]+]] = {{.*}}c"x\00"35__managed__ int x = 1;36 37// DEV-DAG: @v.managed = addrspace(1) externally_initialized global [100 x %struct.vec] zeroinitializer, align 438// DEV-DAG: @v = addrspace(1) externally_initialized global ptr addrspace(1) null39__managed__ vec v[100];40 41// DEV-DAG: @v2.managed = addrspace(1) externally_initialized global <{ %struct.vec, [99 x %struct.vec] }> <{ %struct.vec { float 1.000000e+00, float 1.000000e+00, float 1.000000e+00 }, [99 x %struct.vec] zeroinitializer }>, align 442// DEV-DAG: @v2 = addrspace(1) externally_initialized global ptr addrspace(1) null43__managed__ vec v2[100] = {{1, 1, 1}};44 45// DEV-DAG: @ex.managed = external addrspace(1) global i32, align 446// DEV-DAG: @ex = external addrspace(1) externally_initialized global ptr addrspace(1)47// HOST-DAG: @ex.managed = external global i3248// HOST-DAG: @ex = external externally_initialized global ptr49extern __managed__ int ex;50 51// NORDC-D-DAG: @_ZL2sx.managed = addrspace(1) externally_initialized global i32 1, align 452// NORDC-D-DAG: @_ZL2sx = addrspace(1) externally_initialized global ptr addrspace(1) null53// RDC-D-DAG: @_ZL2sx.static.[[HASH:.*]].managed = addrspace(1) externally_initialized global i32 1, align 454// RDC-D-DAG: @_ZL2sx.static.[[HASH]] = addrspace(1) externally_initialized global ptr addrspace(1) null55// HOST-DAG: @_ZL2sx.managed = internal global i32 156// HOST-DAG: @_ZL2sx = internal externally_initialized global ptr null57// NORDC-DAG: @[[DEVNAMESX:[0-9]+]] = {{.*}}c"_ZL2sx\00"58// RDC-DAG: @[[DEVNAMESX:[0-9]+]] = {{.*}}c"_ZL2sx.static.[[HASH:.*]]\00"59 60// POSTFIX: @_ZL2sx.static.[[HASH:.*]] = addrspace(1) externally_initialized global ptr addrspace(1) null61// POSTFIX: @[[DEVNAMESX:[0-9]+]] = {{.*}}c"_ZL2sx.static.[[HASH]]\00"62static __managed__ int sx = 1;63 64// DEV-DAG: @llvm.compiler.used65// DEV-SAME-DAG: @x.managed66// DEV-SAME-DAG: @x67// DEV-SAME-DAG: @v.managed68// DEV-SAME-DAG: @v69// DEV-SAME-DAG: @_ZL2sx.managed70// DEV-SAME-DAG: @_ZL2sx71 72// Force ex and sx mitted in device compilation.73__global__ void foo(int *z) {74 *z = x + ex + sx;75 v[1].x = 2;76}77 78// Force ex and sx emitted in host compilatioin.79int foo2() {80 return ex + sx;81}82 83// COMMON-LABEL: define {{.*}}@_Z4loadv()84// DEV: %ld.managed = load ptr addrspace(1), ptr addrspace(1) @x, align 485// DEV: %0 = addrspacecast ptr addrspace(1) %ld.managed to ptr86// DEV: %1 = load i32, ptr %0, align 487// DEV: ret i32 %188// HOST: %ld.managed = load ptr, ptr @x, align 489// HOST: %0 = load i32, ptr %ld.managed, align 490// HOST: ret i32 %091__device__ __host__ int load() {92 return x;93}94 95// COMMON-LABEL: define {{.*}}@_Z5storev()96// DEV: %ld.managed = load ptr addrspace(1), ptr addrspace(1) @x, align 497// DEV: %0 = addrspacecast ptr addrspace(1) %ld.managed to ptr98// DEV: store i32 2, ptr %0, align 499// HOST: %ld.managed = load ptr, ptr @x, align 4100// HOST: store i32 2, ptr %ld.managed, align 4101__device__ __host__ void store() {102 x = 2;103}104 105// COMMON-LABEL: define {{.*}}@_Z10addr_takenv()106// DEV: %0 = addrspacecast ptr addrspace(1) %ld.managed to ptr107// DEV: store ptr %0, ptr %p.ascast, align 8108// DEV: %1 = load ptr, ptr %p.ascast, align 8109// DEV: store i32 3, ptr %1, align 4110// HOST: %ld.managed = load ptr, ptr @x, align 4111// HOST: store ptr %ld.managed, ptr %p, align 8112// HOST: %0 = load ptr, ptr %p, align 8113// HOST: store i32 3, ptr %0, align 4114__device__ __host__ void addr_taken() {115 int *p = &x;116 *p = 3;117}118 119// HOST-LABEL: define {{.*}}@_Z5load2v()120// HOST: %ld.managed = load ptr, ptr @v, align 16121// HOST: %0 = getelementptr inbounds [100 x %struct.vec], ptr %ld.managed, i64 0, i64 1122// HOST: %1 = load float, ptr %0, align 4123// HOST: ret float %1124__device__ __host__ float load2() {125 return v[1].x;126}127 128// HOST-LABEL: define {{.*}}@_Z5load3v()129// HOST: %ld.managed = load ptr, ptr @v2, align 16130// HOST: %0 = getelementptr inbounds [100 x %struct.vec], ptr %ld.managed, i64 0, i64 1131// HOST: %1 = getelementptr inbounds nuw %struct.vec, ptr %0, i32 0, i32 1132// HOST: %2 = load float, ptr %1, align 4133// HOST: ret float %2134float load3() {135 return v2[1].y;136}137 138// HOST-LABEL: define {{.*}}@_Z11addr_taken2v()139// HOST: %ld.managed = load ptr, ptr @v, align 16140// HOST: %0 = getelementptr inbounds [100 x %struct.vec], ptr %ld.managed, i64 0, i64 1141// HOST: %1 = ptrtoint ptr %0 to i64142// HOST: %ld.managed1 = load ptr, ptr @v2, align 16143// HOST: %2 = getelementptr inbounds [100 x %struct.vec], ptr %ld.managed1, i64 0, i64 1144// HOST: %3 = getelementptr inbounds nuw %struct.vec, ptr %2, i32 0, i32 1145// HOST: %4 = ptrtoint ptr %3 to i64146// HOST: %5 = sub i64 %4, %1147// HOST: %sub.ptr.div = sdiv exact i64 %5, 4148// HOST: %conv = sitofp i64 %sub.ptr.div to float149// HOST: ret float %conv150float addr_taken2() {151 return (float)reinterpret_cast<long>(&(v2[1].y)-&(v[1].x));152}153 154// COMMON-LABEL: define {{.*}}@_Z5load4v()155// DEV: %ld.managed = load ptr addrspace(1), ptr addrspace(1) @ex, align 4156// DEV: %0 = addrspacecast ptr addrspace(1) %ld.managed to ptr157// DEV: %1 = load i32, ptr %0, align 4158// DEV: ret i32 %1159// HOST: %ld.managed = load ptr, ptr @ex, align 4160// HOST: %0 = load i32, ptr %ld.managed, align 4161// HOST: ret i32 %0162__device__ __host__ int load4() {163 return ex;164}165 166// HOST-DAG: __hipRegisterManagedVar({{.*}}, ptr @x, ptr @x.managed, ptr @[[DEVNAMEX]], i64 4, i32 4)167// HOST-DAG: __hipRegisterManagedVar({{.*}}, ptr @_ZL2sx, ptr @_ZL2sx.managed, ptr @[[DEVNAMESX]]168// HOST-NOT: __hipRegisterManagedVar({{.*}}, ptr @ex, ptr @ex.managed169// HOST-DAG: declare void @__hipRegisterManagedVar(ptr, ptr, ptr, ptr, i64, i32)170