Skip to content

[NFCI][SYCL] Remove srcloc metadata before sycl-post-link #17727

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 4 commits into from
Apr 3, 2025
Merged

[NFCI][SYCL] Remove srcloc metadata before sycl-post-link #17727

Show file tree
Hide file tree
Changes from 1 commit
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
1b344a1
[NFC][SYCL] Remove srcloc metadata before sycl-post-link
MrSidims Mar 29, 2025
a8da38b
apply comments
MrSidims Mar 31, 2025
774587f
fix test, apply comment
MrSidims Mar 31, 2025
1faac6e
Merge remote-tracking branch 'origin/sycl' into clean-up-md
MrSidims Apr 1, 2025
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
19 changes: 16 additions & 3 deletions llvm/lib/SYCLLowerIR/CleanupSYCLMetadata.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -24,8 +24,7 @@
using namespace llvm;

namespace {

void cleanupSYCLCompilerMetadata(const Module &M, llvm::StringRef MD) {
void cleanupSYCLCompilerModuleMetadata(const Module &M, llvm::StringRef MD) {
NamedMDNode *Node = M.getNamedMetadata(MD);
if (!Node)
return;
Expand All @@ -34,6 +33,14 @@ void cleanupSYCLCompilerMetadata(const Module &M, llvm::StringRef MD) {
Node->eraseFromParent();
}

void cleanupSYCLCompilerFunctionMetadata(Function &F,
llvm::StringRef MD) {
MDNode *Node = F.getMetadata(MD);
if (!Node)
return;
F.setMetadata(MD, nullptr);
}

// GV is supposed to be either llvm.compiler.used or llvm.used.
SmallVector<Constant *>
eraseGlobalVariableAndReturnOperands(GlobalVariable *GV) {
Expand Down Expand Up @@ -65,7 +72,13 @@ PreservedAnalyses CleanupSYCLMetadataPass::run(Module &M,
SmallVector<StringRef, 2> ModuleMDToRemove = {"sycl_aspects",
"sycl_types_that_use_aspects"};
for (const auto &MD : ModuleMDToRemove)
cleanupSYCLCompilerMetadata(M, MD);
cleanupSYCLCompilerModuleMetadata(M, MD);

SmallVector<StringRef, 1> FunctionMDToRemove = {"srcloc"};
for (auto &F : M) {
for (const auto &MD : FunctionMDToRemove)
cleanupSYCLCompilerFunctionMetadata(F, MD);
}

return PreservedAnalyses::all();
}
Expand Down
30 changes: 30 additions & 0 deletions llvm/test/SYCLLowerIR/CleanupSYCLCompilerInternalMetadata/sycl-function-metadata.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,30 @@
; RUN: opt -passes=cleanup-sycl-metadata -S < %s | FileCheck %s
;
; Test checks that the pass is able to cleanup sycl_declared_aspects and srcloc
; function metadata

; CHECK-NOT: srcloc

; ModuleID = 'test.cpp'
source_filename = "test.cpp"
target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64-G1"
target triple = "spir64-unknown-unknown"

$_ZNK13KernelFunctorclEv = comdat any

define dso_local spir_func void @_Z6func10v() !sycl_declared_aspects !1 !srcloc !2 {
entry:
ret void
}


define linkonce_odr spir_func void @_ZNK13KernelFunctorclEv() !sycl_declared_aspects !3 !srcloc !4 {
entry:
call spir_func void @_Z6func10v()
ret void
}

!1 = !{i32 5}
!2 = !{i32 2457}
!3 = !{i32 1}
!4 = !{i32 2547}
2 changes: 1 addition & 1 deletion sycl/test/check_device_code/atomic_ref.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,7 +4,7 @@
#include <sycl/sycl.hpp>

// CHECK-LABEL: define dso_local spir_func noundef i32 @_Z17atomic_ref_globalRi(
// CHECK-SAME: ptr addrspace(4) noundef align 4 dereferenceable(4) [[I:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] !srcloc [[META6:![0-9]+]] !sycl_fixed_targets [[META7:![0-9]+]] {
// CHECK-SAME: ptr addrspace(4) noundef align 4 dereferenceable(4) [[I:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] !sycl_fixed_targets [[META6:![0-9]+]] {
// CHECK-NEXT: [[ENTRY:.*:]]
// CHECK-NEXT: [[TMP:%.*]] = addrspacecast ptr addrspace(4) [[I]] to ptr addrspace(1)
// CHECK-NEXT: [[CALL3_I_I:%.*]] = tail call spir_func noundef i32 @_Z18__spirv_AtomicLoadPU3AS1KiN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagE(ptr addrspace(1) noundef [[TMP]], i32 noundef 1, i32 noundef 898) #[[ATTR4:[0-9]+]]
Expand Down
24 changes: 12 additions & 12 deletions sycl/test/check_device_code/bf16_vector_conversion.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -10,7 +10,7 @@ using namespace sycl;
using bfloat16 = sycl::ext::oneapi::bfloat16;

// CHECK-LABEL: define dso_local spir_func void @_Z10TestBFtoF1PN4sycl3_V13ext6oneapi8bfloat16EPf(
// CHECK-SAME: ptr addrspace(4) noundef [[A:%.*]], ptr addrspace(4) noundef [[B:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] !srcloc [[META5:![0-9]+]] !sycl_fixed_targets [[META6:![0-9]+]] {
// CHECK-SAME: ptr addrspace(4) noundef [[A:%.*]], ptr addrspace(4) noundef [[B:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] !sycl_fixed_targets [[META6:![0-9]+]] {
// CHECK-NEXT: entry:
// CHECK-NEXT: tail call spir_func void @__devicelib_ConvertBF16ToFINTELVec1(ptr addrspace(4) noundef [[A]], ptr addrspace(4) noundef [[B]]) #[[ATTR2:[0-9]+]]
// CHECK-NEXT: ret void
Expand All @@ -20,7 +20,7 @@ SYCL_EXTERNAL auto TestBFtoF1(bfloat16 *a, float *b) {
}

// CHECK-LABEL: define dso_local spir_func void @_Z10TestFtoBF1PfPN4sycl3_V13ext6oneapi8bfloat16Ei(
// CHECK-SAME: ptr addrspace(4) noundef [[A:%.*]], ptr addrspace(4) noundef [[B:%.*]], i32 noundef [[SIZE:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META7:![0-9]+]] !sycl_fixed_targets [[META6]] {
// CHECK-SAME: ptr addrspace(4) noundef [[A:%.*]], ptr addrspace(4) noundef [[B:%.*]], i32 noundef [[SIZE:%.*]]) local_unnamed_addr #[[ATTR0]] !sycl_fixed_targets [[META6]] {
// CHECK-NEXT: entry:
// CHECK-NEXT: tail call spir_func void @__devicelib_ConvertFToBF16INTELVec1(ptr addrspace(4) noundef [[A]], ptr addrspace(4) noundef [[B]]) #[[ATTR2]]
// CHECK-NEXT: ret void
Expand All @@ -30,7 +30,7 @@ SYCL_EXTERNAL auto TestFtoBF1(float *a, bfloat16 *b, int size) {
}

// CHECK-LABEL: define dso_local spir_func void @_Z10TestBFtoF2PN4sycl3_V13ext6oneapi8bfloat16EPf(
// CHECK-SAME: ptr addrspace(4) noundef [[A:%.*]], ptr addrspace(4) noundef [[B:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META8:![0-9]+]] !sycl_fixed_targets [[META6]] {
// CHECK-SAME: ptr addrspace(4) noundef [[A:%.*]], ptr addrspace(4) noundef [[B:%.*]]) local_unnamed_addr #[[ATTR0]] !sycl_fixed_targets [[META6]] {
// CHECK-NEXT: entry:
// CHECK-NEXT: tail call spir_func void @__devicelib_ConvertBF16ToFINTELVec2(ptr addrspace(4) noundef [[A]], ptr addrspace(4) noundef [[B]]) #[[ATTR2]]
// CHECK-NEXT: ret void
Expand All @@ -40,7 +40,7 @@ SYCL_EXTERNAL auto TestBFtoF2(bfloat16 *a, float *b) {
}

// CHECK-LABEL: define dso_local spir_func void @_Z10TestFtoBF2PfPN4sycl3_V13ext6oneapi8bfloat16Ei(
// CHECK-SAME: ptr addrspace(4) noundef [[A:%.*]], ptr addrspace(4) noundef [[B:%.*]], i32 noundef [[SIZE:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META9:![0-9]+]] !sycl_fixed_targets [[META6]] {
// CHECK-SAME: ptr addrspace(4) noundef [[A:%.*]], ptr addrspace(4) noundef [[B:%.*]], i32 noundef [[SIZE:%.*]]) local_unnamed_addr #[[ATTR0]] !sycl_fixed_targets [[META6]] {
// CHECK-NEXT: entry:
// CHECK-NEXT: tail call spir_func void @__devicelib_ConvertFToBF16INTELVec2(ptr addrspace(4) noundef [[A]], ptr addrspace(4) noundef [[B]]) #[[ATTR2]]
// CHECK-NEXT: ret void
Expand All @@ -50,7 +50,7 @@ SYCL_EXTERNAL auto TestFtoBF2(float *a, bfloat16 *b, int size) {
}

// CHECK-LABEL: define dso_local spir_func void @_Z10TestBFtoF3PN4sycl3_V13ext6oneapi8bfloat16EPf(
// CHECK-SAME: ptr addrspace(4) noundef [[A:%.*]], ptr addrspace(4) noundef [[B:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META10:![0-9]+]] !sycl_fixed_targets [[META6]] {
// CHECK-SAME: ptr addrspace(4) noundef [[A:%.*]], ptr addrspace(4) noundef [[B:%.*]]) local_unnamed_addr #[[ATTR0]] !sycl_fixed_targets [[META6]] {
// CHECK-NEXT: entry:
// CHECK-NEXT: tail call spir_func void @__devicelib_ConvertBF16ToFINTELVec3(ptr addrspace(4) noundef [[A]], ptr addrspace(4) noundef [[B]]) #[[ATTR2]]
// CHECK-NEXT: ret void
Expand All @@ -60,7 +60,7 @@ SYCL_EXTERNAL auto TestBFtoF3(bfloat16 *a, float *b) {
}

// CHECK-LABEL: define dso_local spir_func void @_Z10TestFtoBF3PfPN4sycl3_V13ext6oneapi8bfloat16Ei(
// CHECK-SAME: ptr addrspace(4) noundef [[A:%.*]], ptr addrspace(4) noundef [[B:%.*]], i32 noundef [[SIZE:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META11:![0-9]+]] !sycl_fixed_targets [[META6]] {
// CHECK-SAME: ptr addrspace(4) noundef [[A:%.*]], ptr addrspace(4) noundef [[B:%.*]], i32 noundef [[SIZE:%.*]]) local_unnamed_addr #[[ATTR0]] !sycl_fixed_targets [[META6]] {
// CHECK-NEXT: entry:
// CHECK-NEXT: tail call spir_func void @__devicelib_ConvertFToBF16INTELVec3(ptr addrspace(4) noundef [[A]], ptr addrspace(4) noundef [[B]]) #[[ATTR2]]
// CHECK-NEXT: ret void
Expand All @@ -70,7 +70,7 @@ SYCL_EXTERNAL auto TestFtoBF3(float *a, bfloat16 *b, int size) {
}

// CHECK-LABEL: define dso_local spir_func void @_Z10TestBFtoF4PN4sycl3_V13ext6oneapi8bfloat16EPf(
// CHECK-SAME: ptr addrspace(4) noundef [[A:%.*]], ptr addrspace(4) noundef [[B:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META12:![0-9]+]] !sycl_fixed_targets [[META6]] {
// CHECK-SAME: ptr addrspace(4) noundef [[A:%.*]], ptr addrspace(4) noundef [[B:%.*]]) local_unnamed_addr #[[ATTR0]] !sycl_fixed_targets [[META6]] {
// CHECK-NEXT: entry:
// CHECK-NEXT: tail call spir_func void @__devicelib_ConvertBF16ToFINTELVec4(ptr addrspace(4) noundef [[A]], ptr addrspace(4) noundef [[B]]) #[[ATTR2]]
// CHECK-NEXT: ret void
Expand All @@ -80,7 +80,7 @@ SYCL_EXTERNAL auto TestBFtoF4(bfloat16 *a, float *b) {
}

// CHECK-LABEL: define dso_local spir_func void @_Z10TestFtoBF4PfPN4sycl3_V13ext6oneapi8bfloat16Ei(
// CHECK-SAME: ptr addrspace(4) noundef [[A:%.*]], ptr addrspace(4) noundef [[B:%.*]], i32 noundef [[SIZE:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META13:![0-9]+]] !sycl_fixed_targets [[META6]] {
// CHECK-SAME: ptr addrspace(4) noundef [[A:%.*]], ptr addrspace(4) noundef [[B:%.*]], i32 noundef [[SIZE:%.*]]) local_unnamed_addr #[[ATTR0]] !sycl_fixed_targets [[META6]] {
// CHECK-NEXT: entry:
// CHECK-NEXT: tail call spir_func void @__devicelib_ConvertFToBF16INTELVec4(ptr addrspace(4) noundef [[A]], ptr addrspace(4) noundef [[B]]) #[[ATTR2]]
// CHECK-NEXT: ret void
Expand All @@ -90,7 +90,7 @@ SYCL_EXTERNAL auto TestFtoBF4(float *a, bfloat16 *b, int size) {
}

// CHECK-LABEL: define dso_local spir_func void @_Z10TestBFtoF8PN4sycl3_V13ext6oneapi8bfloat16EPf(
// CHECK-SAME: ptr addrspace(4) noundef [[A:%.*]], ptr addrspace(4) noundef [[B:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META14:![0-9]+]] !sycl_fixed_targets [[META6]] {
// CHECK-SAME: ptr addrspace(4) noundef [[A:%.*]], ptr addrspace(4) noundef [[B:%.*]]) local_unnamed_addr #[[ATTR0]] !sycl_fixed_targets [[META6]] {
// CHECK-NEXT: entry:
// CHECK-NEXT: tail call spir_func void @__devicelib_ConvertBF16ToFINTELVec8(ptr addrspace(4) noundef [[A]], ptr addrspace(4) noundef [[B]]) #[[ATTR2]]
// CHECK-NEXT: ret void
Expand All @@ -100,7 +100,7 @@ SYCL_EXTERNAL auto TestBFtoF8(bfloat16 *a, float *b) {
}

// CHECK-LABEL: define dso_local spir_func void @_Z10TestFtoBF8PfPN4sycl3_V13ext6oneapi8bfloat16Ei(
// CHECK-SAME: ptr addrspace(4) noundef [[A:%.*]], ptr addrspace(4) noundef [[B:%.*]], i32 noundef [[SIZE:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META15:![0-9]+]] !sycl_fixed_targets [[META6]] {
// CHECK-SAME: ptr addrspace(4) noundef [[A:%.*]], ptr addrspace(4) noundef [[B:%.*]], i32 noundef [[SIZE:%.*]]) local_unnamed_addr #[[ATTR0]] !sycl_fixed_targets [[META6]] {
// CHECK-NEXT: entry:
// CHECK-NEXT: tail call spir_func void @__devicelib_ConvertFToBF16INTELVec8(ptr addrspace(4) noundef [[A]], ptr addrspace(4) noundef [[B]]) #[[ATTR2]]
// CHECK-NEXT: ret void
Expand All @@ -110,7 +110,7 @@ SYCL_EXTERNAL auto TestFtoBF8(float *a, bfloat16 *b, int size) {
}

// CHECK-LABEL: define dso_local spir_func void @_Z11TestBFtoF16PN4sycl3_V13ext6oneapi8bfloat16EPf(
// CHECK-SAME: ptr addrspace(4) noundef [[A:%.*]], ptr addrspace(4) noundef [[B:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META16:![0-9]+]] !sycl_fixed_targets [[META6]] {
// CHECK-SAME: ptr addrspace(4) noundef [[A:%.*]], ptr addrspace(4) noundef [[B:%.*]]) local_unnamed_addr #[[ATTR0]] !sycl_fixed_targets [[META6]] {
// CHECK-NEXT: entry:
// CHECK-NEXT: tail call spir_func void @__devicelib_ConvertBF16ToFINTELVec16(ptr addrspace(4) noundef [[A]], ptr addrspace(4) noundef [[B]]) #[[ATTR2]]
// CHECK-NEXT: ret void
Expand All @@ -120,7 +120,7 @@ SYCL_EXTERNAL auto TestBFtoF16(bfloat16 *a, float *b) {
}

// CHECK-LABEL: define dso_local spir_func void @_Z11TestFtoBF16PfPN4sycl3_V13ext6oneapi8bfloat16Ei(
// CHECK-SAME: ptr addrspace(4) noundef [[A:%.*]], ptr addrspace(4) noundef [[B:%.*]], i32 noundef [[SIZE:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META17:![0-9]+]] !sycl_fixed_targets [[META6]] {
// CHECK-SAME: ptr addrspace(4) noundef [[A:%.*]], ptr addrspace(4) noundef [[B:%.*]], i32 noundef [[SIZE:%.*]]) local_unnamed_addr #[[ATTR0]] !sycl_fixed_targets [[META6]] {
// CHECK-NEXT: entry:
// CHECK-NEXT: tail call spir_func void @__devicelib_ConvertFToBF16INTELVec16(ptr addrspace(4) noundef [[A]], ptr addrspace(4) noundef [[B]]) #[[ATTR2]]
// CHECK-NEXT: ret void
Expand Down
17 changes: 0 additions & 17 deletions sycl/test/check_device_code/device_has_func.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,6 +1,5 @@
// RUN: %clangxx -fsycl -Xclang -fsycl-is-device -fsycl-device-only -Xclang -fno-sycl-early-optimizations -S -emit-llvm %s -o %t.ll
// RUN: FileCheck %s --input-file %t.ll --check-prefix=CHECK-ASPECTS
// RUN: FileCheck %s --input-file %t.ll --check-prefix=CHECK-SRCLOC

// Tests for IR of device_has(aspect, ...) attribute and
// !sycl_used_aspects metadata.
Expand All @@ -11,37 +10,30 @@
using namespace sycl;

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

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

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

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

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

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

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

SYCL_EXTERNAL [[sycl::device_has(sycl::aspect::cpu)]] void kernel_name_1() {
Expand All @@ -54,23 +46,14 @@ SYCL_EXTERNAL [[sycl::device_has(sycl::aspect::cpu)]] void kernel_name_1() {
}

// CHECK-ASPECTS: define dso_local spir_func void @{{.*}}kernel_name_2{{.*}} !sycl_declared_aspects ![[ASPECTS4:[0-9]+]]
// CHECK-SRCLOC: define dso_local spir_func void @{{.*}}kernel_name_2{{.*}} !srcloc ![[SRCLOC8:[0-9]+]] {{.*}}
SYCL_EXTERNAL [[sycl::device_has(sycl::aspect::gpu)]] void kernel_name_2() {}

// CHECK-ASPECTS-DAG: [[ASPECTS1]] = !{![[ASPECTCPU:[0-9]+]]}
// CHECK-ASPECTS-DAG: [[ASPECTCPU]] = !{!"cpu", i32 1}
// CHECK-SRCLOC-DAG: [[SRCLOC1]] = !{i32 {{[0-9]+}}}
// CHECK-ASPECTS-DAG: [[EMPTYASPECTS]] = !{}
// CHECK-SRCLOC-DAG: [[SRCLOC2]] = !{i32 {{[0-9]+}}}
// CHECK-ASPECTS-DAG: [[ASPECTS2]] = !{![[ASPECTFP16:[0-9]+]], ![[ASPECTGPU:[0-9]+]]}
// CHECK-ASPECTS-DAG: [[ASPECTFP16]] = !{!"fp16", i32 5}
// CHECK-ASPECTS-DAG: [[ASPECTGPU]] = !{!"gpu", i32 2}
// CHECK-SRCLOC-DAG: [[SRCLOC3]] = !{i32 {{[0-9]+}}}
// CHECK-SRCLOC-DAG: [[SRCLOC4]] = !{i32 {{[0-9]+}}}
// CHECK-ASPECTS-DAG: [[ASPECTS3]] = !{![[ASPECTHOST:[0-9]+]]}
// CHECK-ASPECTS-DAG: [[ASPECTHOST]] = !{!"host", i32 0}
// CHECK-SRCLOC-DAG: [[SRCLOC5]] = !{i32 {{[0-9]+}}}
// CHECK-SRCLOC-DAG: [[SRCLOC6]] = !{i32 {{[0-9]+}}}
// CHECK-SRCLOC-DAG: [[SRCLOC7]] = !{i32 {{[0-9]+}}}
// CHECK-ASPECTS-DAG: [[ASPECTS4]] = !{![[ASPECTGPU]]}
// CHECK-SRCLOC-DAG: [[SRCLOC8]] = !{i32 {{[0-9]+}}}
17 changes: 0 additions & 17 deletions sycl/test/check_device_code/device_has_kernel.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,6 +1,5 @@
// RUN: %clangxx -fsycl -Xclang -fsycl-is-device -fsycl-device-only -Xclang -fno-sycl-early-optimizations -S -emit-llvm %s -o %t.ll
// RUN: FileCheck %s --input-file %t.ll --check-prefix=CHECK-ASPECTS
// RUN: FileCheck %s --input-file %t.ll --check-prefix=CHECK-SRCLOC

// Tests for IR of device_has(aspect, ...) attribute and
// !sycl_used_aspects metadata.
Expand All @@ -12,37 +11,30 @@ using namespace sycl;
queue q;

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

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

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

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

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

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

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

class KernelFunctor {
Expand All @@ -62,26 +54,17 @@ void foo() {
KernelFunctor f1;
h.single_task<class kernel_name_1>(f1);
// CHECK-ASPECTS: define weak_odr dso_local spir_kernel void @{{.*}}kernel_name_2{{.*}} !sycl_declared_aspects ![[ASPECTS4:[0-9]+]]
// CHECK-SRCLOC: define weak_odr dso_local spir_kernel void @{{.*}}kernel_name_2{{.*}} !srcloc ![[SRCLOC8:[0-9]+]] {{.*}}
h.single_task<class kernel_name_2>(
[]() [[sycl::device_has(sycl::aspect::gpu)]] {});
});
}

// CHECK-ASPECTS-DAG: [[ASPECTS1]] = !{![[ASPECTCPU:[0-9]+]]}
// CHECK-ASPECTS-DAG: [[ASPECTCPU]] = !{!"cpu", i32 1}
// CHECK-SRCLOC-DAG: [[SRCLOC1]] = !{i32 {{[0-9]+}}}
// CHECK-ASPECTS-DAG: [[EMPTYASPECTS]] = !{}
// CHECK-SRCLOC-DAG: [[SRCLOC2]] = !{i32 {{[0-9]+}}}
// CHECK-ASPECTS-DAG: [[ASPECTS2]] = !{![[ASPECTFP16:[0-9]+]], ![[ASPECTGPU:[0-9]+]]}
// CHECK-ASPECTS-DAG: [[ASPECTFP16]] = !{!"fp16", i32 5}
// CHECK-ASPECTS-DAG: [[ASPECTGPU]] = !{!"gpu", i32 2}
// CHECK-SRCLOC-DAG: [[SRCLOC3]] = !{i32 {{[0-9]+}}}
// CHECK-SRCLOC-DAG: [[SRCLOC4]] = !{i32 {{[0-9]+}}}
// CHECK-ASPECTS-DAG: [[ASPECTS3]] = !{![[ASPECTHOST:[0-9]+]]}
// CHECK-ASPECTS-DAG: [[ASPECTHOST]] = !{!"host", i32 0}
// CHECK-SRCLOC-DAG: [[SRCLOC5]] = !{i32 {{[0-9]+}}}
// CHECK-SRCLOC-DAG: [[SRCLOC6]] = !{i32 {{[0-9]+}}}
// CHECK-SRCLOC-DAG: [[SRCLOC7]] = !{i32 {{[0-9]+}}}
// CHECK-ASPECTS-DAG: [[ASPECTS4]] = !{![[ASPECTGPU]]}
// CHECK-SRCLOC-DAG: [[SRCLOC8]] = !{i32 {{[0-9]+}}}
Loading
Loading