Skip to content

[SYCL] Change linkage to linkonce_odr unless SYCL_EXTERNAL. #10317

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 1 commit into from
Sep 12, 2023
Merged

[SYCL] Change linkage to linkonce_odr unless SYCL_EXTERNAL. #10317

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 1 addition & 1 deletion clang/include/clang/Sema/Sema.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -14159,7 +14159,7 @@ class Sema final {
SourceLocation RParenLoc);

template <typename AttrTy>
bool isTypeDecoratedWithDeclAttribute(QualType Ty) {
static bool isTypeDecoratedWithDeclAttribute(QualType Ty) {
const CXXRecordDecl *RecTy = Ty->getAsCXXRecordDecl();
if (!RecTy)
return false;
Expand Down
14 changes: 14 additions & 0 deletions clang/lib/CodeGen/CodeGenModule.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -48,6 +48,7 @@
#include "clang/CodeGen/BackendUtil.h"
#include "clang/CodeGen/ConstantInitBuilder.h"
#include "clang/Frontend/FrontendDiagnostic.h"
#include "clang/Sema/Sema.h"
#include "llvm/ADT/STLExtras.h"
#include "llvm/ADT/StringExtras.h"
#include "llvm/ADT/StringSwitch.h"
Expand Down Expand Up @@ -6082,6 +6083,19 @@ llvm::GlobalValue::LinkageTypes CodeGenModule::getLLVMLinkageForDeclarator(
if (Linkage == GVA_AvailableExternally)
return llvm::GlobalValue::AvailableExternallyLinkage;

// SYCL: Device code is not generally limited to one translation unit, but
// anything accessed from another translation unit is required to be annotated
// with the SYCL_EXTERNAL macro. For any function or variable that does not
// have this, linkonce_odr suffices. If -fno-sycl-rdc is passed, we know there
// is only one translation unit and can so mark them internal.
if (getLangOpts().SYCLIsDevice && !D->hasAttr<SYCLKernelAttr>() &&
!D->hasAttr<SYCLDeviceAttr>() &&
!Sema::isTypeDecoratedWithDeclAttribute<SYCLDeviceGlobalAttr>(
D->getType()))
return getLangOpts().GPURelocatableDeviceCode
? llvm::Function::LinkOnceODRLinkage
: llvm::Function::InternalLinkage;

// Note that Apple's kernel linker doesn't support symbol
// coalescing, so we need to avoid linkonce and weak linkages there.
// Normally, this means we just map to internal, but for explicit
Expand Down
2 changes: 1 addition & 1 deletion clang/test/CodeGenSYCL/Inputs/sycl.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -6,7 +6,7 @@
extern "C" int printf(const char* fmt, ...);

#ifdef __SYCL_DEVICE_ONLY__
__attribute__((convergent)) extern SYCL_EXTERNAL void
__attribute__((convergent)) extern __attribute__((sycl_device)) void
__spirv_ControlBarrier(int, int, int) noexcept;
#endif

Expand Down
18 changes: 9 additions & 9 deletions clang/test/CodeGenSYCL/check-direct-attribute-propagation.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -175,7 +175,7 @@ int main() {
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4() #0{{.*}} !kernel_arg_buffer_location ![[NUM]]
// CHECK-NOT: !scheduler_target_fmax_mhz
// CHECK-SAME: {
// CHECK: define dso_local spir_func void @_Z3foov()
// CHECK: define {{.*}}spir_func void @_Z3foov()
h.single_task<class kernel_name4>(
[]() { foo(); });

Expand All @@ -195,7 +195,7 @@ int main() {
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name8() #0{{.*}} !kernel_arg_buffer_location ![[NUM]]
// CHECK-NOT: !num_simd_work_items
// CHECK-SAME: {
// CHECK: define dso_local spir_func void @_Z4foo1v()
// CHECK: define {{.*}}spir_func void @_Z4foo1v()
h.single_task<class kernel_name8>(
[]() { foo1(); });

Expand All @@ -215,7 +215,7 @@ int main() {
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name12() #0{{.*}} !kernel_arg_buffer_location ![[NUM]]
// CHECK-NOT: !no_global_work_offset
// CHECK-SAME: {
// CHECK: define dso_local spir_func void @_Z4foo2v()
// CHECK: define {{.*}}spir_func void @_Z4foo2v()
h.single_task<class kernel_name12>(
[]() { foo2(); });

Expand All @@ -235,7 +235,7 @@ int main() {
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name16() #0{{.*}} !kernel_arg_buffer_location ![[NUM]]
// CHECK-NOT: !max_global_work_dim
// CHECK-SAME: {
// CHECK: define dso_local spir_func void @_Z4foo3v()
// CHECK: define {{.*}}spir_func void @_Z4foo3v()
h.single_task<class kernel_name16>(
[]() { foo3(); });

Expand All @@ -255,7 +255,7 @@ int main() {
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name20() #0{{.*}} !kernel_arg_buffer_location ![[NUM]]
// CHECK-NOT: !reqd_sub_group_size
// CHECK-SAME: {
// CHECK: define dso_local spir_func void @_Z4foo4v()
// CHECK: define {{.*}}spir_func void @_Z4foo4v()
Functor4 f4;
h.single_task<class kernel_name20>(f4);

Expand All @@ -275,7 +275,7 @@ int main() {
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name24() #0{{.*}} !kernel_arg_buffer_location ![[NUM]]
// CHECK-NOT: !reqd_work_group_size
// CHECK-SAME: {
// CHECK: define dso_local spir_func void @_Z4foo5v()
// CHECK: define {{.*}}spir_func void @_Z4foo5v()
Functor6 f6;
h.single_task<class kernel_name24>(f6);

Expand All @@ -295,7 +295,7 @@ int main() {
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name28() #0{{.*}} !kernel_arg_buffer_location ![[NUM]]
// CHECK-NOT: !max_work_group_size
// CHECK-SAME: {
// CHECK: define dso_local spir_func void @_Z4foo6v()
// CHECK: define {{.*}}spir_func void @_Z4foo6v()
Functor8 f8;
h.single_task<class kernel_name28>(f8);

Expand All @@ -320,7 +320,7 @@ int main() {
// CHECK: define {{.*}}spir_func void @{{.*}}Functor10{{.*}}(ptr addrspace(4) noundef align 1 dereferenceable_or_null(1) %this) #3 comdat align 2
// CHECK-NOT: noalias
// CHECK-SAME: {
// CHECK: define dso_local spir_func void @_Z4foo8v()
// CHECK: define {{.*}}spir_func void @_Z4foo8v()
Functor10 f10;
h.single_task<class kernel_name32>(f10);

Expand Down Expand Up @@ -350,7 +350,7 @@ int main() {
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name38()
// CHECK-NOT: !work_group_size_hint
// CHECK-SAME: {
// CHECK: define dso_local spir_func void @_Z5foo11v()
// CHECK: define {{.*}}spir_func void @_Z5foo11v()
h.single_task<class kernel_name38>(
[]() { foo11(); });

Expand Down
12 changes: 8 additions & 4 deletions clang/test/CodeGenSYCL/device_global.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,4 +1,5 @@
// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -fsycl-unique-prefix=THE_PREFIX -std=c++17 -emit-llvm %s -o - | FileCheck %s
// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -fsycl-unique-prefix=THE_PREFIX -std=c++17 -fgpu-rdc -emit-llvm %s -o - | FileCheck %s --check-prefixes=CHECK,CHECK-RDC
// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -fsycl-unique-prefix=THE_PREFIX -std=c++17 -fno-gpu-rdc -emit-llvm %s -o - | FileCheck %s --check-prefixes=CHECK,CHECK-NORDC
#include "sycl.hpp"

// Test cases below show that 'sycl-unique-id' LLVM IR attribute is attached to the
Expand All @@ -10,14 +11,16 @@ using namespace sycl;
queue q;

device_global<int> A;
#ifdef SYCL_EXTERNAL
SYCL_EXTERNAL device_global<int> AExt;
#endif
static device_global<int> B;

struct Foo {
static device_global<int> C;
};
device_global<int> Foo::C;
// CHECK: @AExt = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global" zeroinitializer, align 8 #[[AEXT_ATTRS:[0-9]+]]
// CHECK-RDC: @AExt = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global" zeroinitializer, align 8 #[[AEXT_ATTRS:[0-9]+]]
// CHECK: @A = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global" zeroinitializer, align 8 #[[A_ATTRS:[0-9]+]]
// CHECK: @_ZL1B = internal addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global" zeroinitializer, align 8 #[[B_ATTRS:[0-9]+]]
// CHECK: @_ZN3Foo1CE = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global" zeroinitializer, align 8 #[[C_ATTRS:[0-9]+]]
Expand All @@ -44,7 +47,8 @@ class [[__sycl_detail__::global_variable_allowed]] only_global_var_allowed {
// check that we don't generate `sycl-unique-id` IR attribute if class does not use
// [[__sycl_detail__::device_global]]
only_global_var_allowed<int> no_device_global;
// CHECK: @no_device_global = addrspace(1) global %class.only_global_var_allowed zeroinitializer, align 8{{$}}
// CHECK-RDC: @no_device_global = linkonce_odr addrspace(1) global %class.only_global_var_allowed zeroinitializer, align 8{{$}}
// CHECK-NORDC: @no_device_global = internal addrspace(1) global %class.only_global_var_allowed zeroinitializer, align 8{{$}}

inline namespace Bar {
device_global<float> InlineNS;
Expand Down Expand Up @@ -97,7 +101,7 @@ void bar() {
// CHECK-SAME: @_ZL1B
// CHECK-SAME: @_ZN12_GLOBAL__N_19same_nameE

// CHECK: attributes #[[AEXT_ATTRS]] = { "sycl-unique-id"="_Z4AExt" }
// CHECK-RDC: attributes #[[AEXT_ATTRS]] = { "sycl-unique-id"="_Z4AExt" }
// CHECK: attributes #[[A_ATTRS]] = { "sycl-unique-id"="_Z1A" }
// CHECK: attributes #[[B_ATTRS]] = { "sycl-unique-id"="THE_PREFIX____ZL1B" }
// CHECK: attributes #[[C_ATTRS]] = { "sycl-unique-id"="_ZN3Foo1CE" }
Expand Down
18 changes: 9 additions & 9 deletions clang/test/CodeGenSYCL/device_has.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -8,33 +8,33 @@ queue q;

// CHECK-DAG: define dso_local spir_kernel void @{{.*}}kernel_name_1{{.*}} !sycl_declared_aspects ![[ASPECTS1:[0-9]+]] !srcloc ![[SRCLOC1:[0-9]+]]

// CHECK-DAG: define dso_local spir_func void @{{.*}}func1{{.*}} !sycl_declared_aspects ![[ASPECTS1]] !srcloc ![[SRCLOC2:[0-9]+]] {
// CHECK-DAG: define {{.*}}spir_func void @{{.*}}func1{{.*}} !sycl_declared_aspects ![[ASPECTS1]] !srcloc ![[SRCLOC2:[0-9]+]] {
[[sycl::device_has(sycl::aspect::cpu)]] void func1() {}

// CHECK-DAG: define dso_local spir_func void @{{.*}}func2{{.*}} !sycl_declared_aspects ![[ASPECTS2:[0-9]+]] !srcloc ![[SRCLOC3:[0-9]+]] {
// CHECK-DAG: define {{.*}}spir_func void @{{.*}}func2{{.*}} !sycl_declared_aspects ![[ASPECTS2:[0-9]+]] !srcloc ![[SRCLOC3:[0-9]+]] {
[[sycl::device_has(sycl::aspect::fp16, sycl::aspect::gpu)]] void func2() {}

// CHECK-DAG: define dso_local spir_func void @{{.*}}func3{{.*}} !sycl_declared_aspects ![[EMPTYASPECTS:[0-9]+]] !srcloc ![[SRCLOC4:[0-9]+]] {
// CHECK-DAG: define {{.*}}spir_func void @{{.*}}func3{{.*}} !sycl_declared_aspects ![[EMPTYASPECTS:[0-9]+]] !srcloc ![[SRCLOC4:[0-9]+]] {
[[sycl::device_has()]] void func3() {}

// CHECK-DAG: define linkonce_odr spir_func void @{{.*}}func4{{.*}} !sycl_declared_aspects ![[ASPECTS3:[0-9]+]] !srcloc ![[SRCLOC5:[0-9]+]] {
// CHECK-DAG: define {{.*}}spir_func void @{{.*}}func4{{.*}} !sycl_declared_aspects ![[ASPECTS3:[0-9]+]] !srcloc ![[SRCLOC5:[0-9]+]] {
template <sycl::aspect Aspect>
[[sycl::device_has(Aspect)]] void func4() {}

// CHECK-DAG: define dso_local spir_func void @{{.*}}func5{{.*}} !sycl_declared_aspects ![[ASPECTS1]] !srcloc ![[SRCLOC6:[0-9]+]] {
// CHECK-DAG: define {{.*}}spir_func void @{{.*}}func5{{.*}} !sycl_declared_aspects ![[ASPECTS1]] !srcloc ![[SRCLOC6:[0-9]+]] {
[[sycl::device_has(sycl::aspect::cpu)]] void func5();
void func5() {}

constexpr sycl::aspect getAspect() { return sycl::aspect::cpu; }
// CHECK-DAG: define dso_local spir_func void @{{.*}}func6{{.*}} !sycl_declared_aspects ![[ASPECTS1]] !srcloc ![[SRCLOC7:[0-9]+]] {
// CHECK-DAG: define {{.*}}spir_func void @{{.*}}func6{{.*}} !sycl_declared_aspects ![[ASPECTS1]] !srcloc ![[SRCLOC7:[0-9]+]] {
[[sycl::device_has(getAspect())]] void func6() {}

// CHECK-DAG: define linkonce_odr spir_func void @{{.*}}func7{{.*}} !sycl_declared_aspects ![[ASPECTS1]]
// CHECK-DAG: define linkonce_odr spir_func void @{{.*}}func7{{.*}} !sycl_declared_aspects ![[ASPECTS5:[0-9]+]]
// CHECK-DAG: define {{.*}}spir_func void @{{.*}}func7{{.*}} !sycl_declared_aspects ![[ASPECTS1]]
// CHECK-DAG: define {{.*}}spir_func void @{{.*}}func7{{.*}} !sycl_declared_aspects ![[ASPECTS5:[0-9]+]]
template <sycl::aspect... Asp>
[[sycl::device_has(Asp...)]] void func7() {}

// CHECK-DAG: define linkonce_odr spir_func void @{{.*}}func8{{.*}} !sycl_declared_aspects ![[ASPECTS5]]
// CHECK-DAG: define {{.*}}spir_func void @{{.*}}func8{{.*}} !sycl_declared_aspects ![[ASPECTS5]]
template <sycl::aspect Asp, sycl::aspect... AspPack>
[[sycl::device_has(Asp, AspPack...)]] void func8() {}

Expand Down
2 changes: 1 addition & 1 deletion clang/test/CodeGenSYCL/function-attrs.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3,7 +3,7 @@

int foo();

// CHECK: define dso_local spir_func void @_Z3barv() [[BAR:#[0-9]+]]
// CHECK: define {{.*}}spir_func void @_Z3barv() [[BAR:#[0-9]+]]
// CHECK: attributes [[BAR]] =
// CHECK-SAME: convergent
// CHECK-SAME: nounwind
Expand Down
2 changes: 1 addition & 1 deletion clang/test/CodeGenSYCL/functionptr-addrspace.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -7,7 +7,7 @@ __attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) {
kernelFunc();
}

// CHECK: define dso_local spir_func{{.*}}invoke_function{{.*}}(ptr noundef %fptr, ptr addrspace(4) noundef %ptr)
// CHECK: define {{.*}}spir_func{{.*}}invoke_function{{.*}}(ptr noundef %fptr, ptr addrspace(4) noundef %ptr)
void invoke_function(int (*fptr)(), int *ptr) {}

int f() { return 0; }
Expand Down
6 changes: 3 additions & 3 deletions clang/test/CodeGenSYCL/loop_unroll.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,7 +1,7 @@
// RUN: %clang_cc1 -triple spir64-unknown-unknown -disable-llvm-passes -fsycl-is-device -emit-llvm %s -o - | FileCheck %s

void enable() {
// CHECK-LABEL: define dso_local spir_func void @_Z6enablev()
// CHECK-LABEL: define {{.*}}spir_func void @_Z6enablev()
int i = 1000;
// CHECK: br i1 %{{.*}}, label %do.body, label %do.end, !llvm.loop ![[ENABLE:[0-9]+]]
[[clang::loop_unroll]]
Expand All @@ -10,7 +10,7 @@ void enable() {

template <int A>
void count() {
// CHECK-LABEL: define linkonce_odr spir_func void @_Z5countILi4EEvv()
// CHECK-LABEL: define {{.*}}spir_func void @_Z5countILi4EEvv()
// CHECK: br label %for.cond, !llvm.loop ![[COUNT:[0-9]+]]
[[clang::loop_unroll(8)]]
for (int i = 0; i < 1000; ++i);
Expand All @@ -21,7 +21,7 @@ void count() {

template <int A>
void disable() {
// CHECK-LABEL: define linkonce_odr spir_func void @_Z7disableILi1EEvv()
// CHECK-LABEL: define {{.*}}spir_func void @_Z7disableILi1EEvv()
int i = 1000, j = 100;
// CHECK: br label %while.cond, !llvm.loop ![[DISABLE:[0-9]+]]
[[clang::loop_unroll(1)]]
Expand Down
2 changes: 1 addition & 1 deletion clang/test/CodeGenSYCL/nontrivial_device_copyable.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -29,7 +29,7 @@ int main() {

// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name(ptr noundef byval(%struct.NontriviallyCopyable)
// CHECK-NOT: define {{.*}}spir_func void @{{.*}}device_func{{.*}}({{.*}}byval(%struct.NontriviallyCopyable)
// CHECK: define dso_local spir_func void @_Z11device_func20NontriviallyCopyable(ptr noundef %X)
// CHECK: define {{.*}}spir_func void @_Z11device_func20NontriviallyCopyable(ptr noundef %X)
// CHECK: %X.indirect_addr = alloca ptr addrspace(4)
// CHECK: %X.indirect_addr.ascast = addrspacecast ptr %X.indirect_addr to ptr addrspace(4)
// CHECK: %X.ascast = addrspacecast ptr %X to ptr addrspace(4)
Expand Down
7 changes: 4 additions & 3 deletions clang/test/CodeGenSYCL/sycl-cuda-host-device-functions.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,6 @@
// RUN: %clang_cc1 -fsycl-is-host -sycl-std=2020 -emit-llvm %s -o - | FileCheck %s -check-prefix CHECK-HOST
// RUN: %clang_cc1 -fsycl-is-device -sycl-std=2020 -emit-llvm %s -o - | FileCheck %s -check-prefix CHECK-DEV
// RUN: %clang_cc1 -fsycl-is-device -sycl-std=2020 -fno-gpu-rdc -emit-llvm %s -o - | FileCheck %s -check-prefixes CHECK-DEV,CHECK-DEV-NORDC
// RUN: %clang_cc1 -fsycl-is-device -sycl-std=2020 -fgpu-rdc -emit-llvm %s -o - | FileCheck %s -check-prefixes CHECK-DEV,CHECK-DEV-RDC

// This tests
// - if a dummy __host__ function (returning undef) is generated for every
Expand Down Expand Up @@ -63,7 +64,8 @@ __device__ int fun5() { return 5; }

int fun6() { return 7; }

// CHECK-DEV: define dso_local noundef i32 @{{.*}}fun6{{.*}}()
// CHECK-DEV-RDC: define linkonce_odr{{ dso_local | }}noundef i32 @{{.*}}fun6{{.*}}()
// CHECK-DEV-NORDC: define internal noundef i32 @{{.*}}fun6{{.*}}()
// CHECK-DEV: ret i32 7

__attribute((sycl_device)) void test() {
Expand All @@ -75,4 +77,3 @@ __attribute((sycl_device)) void test() {
fun5();
fun6();
}

8 changes: 4 additions & 4 deletions clang/test/CodeGenSYCL/sycl-device-static-init.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,11 +1,11 @@
// RUN: %clang_cc1 -fno-sycl-force-inline-kernel-lambda -fsycl-is-device -triple spir64-unknown-unknown -disable-llvm-passes %s -emit-llvm -o - | FileCheck %s
// Test that static initializers do not force the emission of globals on sycl device

// CHECK-NOT: $_ZN8BaseInitI12TestBaseTypeE15s_regbase_ncsdmE = comdat any
// CHECK-NOT: $_ZN8BaseInitI12TestBaseTypeE15s_regbase_ncsdmE =
// CHECK: $_ZN8BaseInitI12TestBaseTypeE3varE = comdat any
// CHECK: @_ZN8BaseInitI12TestBaseTypeE3varE = weak_odr addrspace(1) constant i32 9, comdat, align 4
// CHECK-NOT: @_ZN8BaseInitI12TestBaseTypeE15s_regbase_ncsdmE = weak_odr addrspace(1) global %struct._ZTS16RegisterBaseInit.RegisterBaseInit zeroinitializer, comdat, align 1
// CHECK-NOT: @_ZGVN8BaseInitI12TestBaseTypeE15s_regbase_ncsdmE = weak_odr global i64 0, comdat($_ZN8BaseInitI12TestBaseTypeE9s_regbaseE), align 8
// CHECK: @_ZN8BaseInitI12TestBaseTypeE3varE = {{.*}}addrspace(1) constant i32 9, comdat, align 4
// CHECK-NOT: @_ZN8BaseInitI12TestBaseTypeE15s_regbase_ncsdmE =
// CHECK-NOT: @_ZGVN8BaseInitI12TestBaseTypeE15s_regbase_ncsdmE =
// CHECK: define {{.*}}spir_kernel void @_ZTSZ4mainE11fake_kernel()
// CHECK: call spir_func void @_ZZ4mainENKUlvE_clEv

Expand Down
23 changes: 15 additions & 8 deletions clang/test/CodeGenSYCL/uses_aspects.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -10,35 +10,42 @@ class [[__sycl_detail__::__uses_aspects__(sycl::aspect::cpu)]] Type1WithAspect{}
class [[__sycl_detail__::__uses_aspects__(sycl::aspect::fp16, sycl::aspect::cpu)]] Type2WithAspect{};
class [[__sycl_detail__::__uses_aspects__(sycl::aspect::host)]] UnusedType3WithAspect{};

// CHECK: define dso_local spir_func void @{{.*}}func1{{.*}} !sycl_used_aspects ![[ASPECTS1:[0-9]+]] {
// CHECK: define {{.*}}spir_func void @{{.*}}func1
// CHECK-SAME: !sycl_used_aspects ![[ASPECTS1:[0-9]+]]
[[__sycl_detail__::__uses_aspects__(sycl::aspect::cpu)]] void func1() {}

// CHECK: define dso_local spir_func void @{{.*}}func2{{.*}} !sycl_used_aspects ![[ASPECTS2:[0-9]+]] {
// CHECK: define {{.*}}spir_func void @{{.*}}func2
// CHECK-SAME: !sycl_used_aspects ![[ASPECTS2:[0-9]+]]
[[__sycl_detail__::__uses_aspects__(sycl::aspect::fp16, sycl::aspect::gpu)]] void func2() {}

// CHECK: define dso_local spir_func void @{{.*}}func3{{.*}} !sycl_used_aspects ![[EMPTYASPECTS:[0-9]+]] {
// CHECK: define {{.*}}spir_func void @{{.*}}func3
// CHECK-SAME: !sycl_used_aspects ![[EMPTYASPECTS:[0-9]+]]
[[__sycl_detail__::__uses_aspects__()]] void func3() {}

// CHECK: define linkonce_odr spir_func void @{{.*}}func4{{.*}} !sycl_used_aspects ![[ASPECTS3:[0-9]+]] {
// CHECK: define {{.*}}spir_func void @{{.*}}func4
// CHECK-SAME: !sycl_used_aspects ![[ASPECTS3:[0-9]+]]
template <sycl::aspect Aspect>
[[__sycl_detail__::__uses_aspects__(Aspect)]] void func4() {}

// CHECK: define dso_local spir_func void @{{.*}}func5{{.*}} !sycl_used_aspects ![[ASPECTS1]] {
// CHECK: define {{.*}}spir_func void @{{.*}}func5
// CHECK-SAME: !sycl_used_aspects ![[ASPECTS1]]
[[__sycl_detail__::__uses_aspects__(sycl::aspect::cpu)]] void func5();
void func5() {}

// CHECK: define {{.*}}spir_func void @{{.*}}func6
// CHECK-SAME: !sycl_used_aspects ![[ASPECTS1]]
[[__sycl_detail__::__uses_aspects__(sycl::aspect::cpu)]] void func6();
// CHECK: define dso_local spir_func void @{{.*}}func6{{.*}} !sycl_used_aspects ![[ASPECTS1]] {
void func6() {
Type1WithAspect TestObj1;
Type2WithAspect TestObj2;
}

constexpr sycl::aspect getAspect() { return sycl::aspect::cpu; }
// CHECK: define dso_local spir_func void @{{.*}}func7{{.*}} !sycl_used_aspects ![[ASPECTS1]] {
// CHECK: define {{.*}}spir_func void @{{.*}}func7
// CHECK-SAME: !sycl_used_aspects ![[ASPECTS1]]
[[__sycl_detail__::__uses_aspects__(getAspect())]] void func7() {}

// CHECK: declare !sycl_used_aspects ![[ASPECTS1]] spir_func void @{{.*}}func8{{.*}}
// CHECK: declare !sycl_used_aspects ![[ASPECTS1]] spir_func void @{{.*}}func8
[[__sycl_detail__::__uses_aspects__(sycl::aspect::cpu)]] SYCL_EXTERNAL void func8();

class KernelFunctor {
Expand Down
Loading