Skip to content
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
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
4 changes: 2 additions & 2 deletions clang/test/CodeGenSYCL/sub-group-size.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -15,15 +15,15 @@ using namespace sycl;
// ALL-DAG: define {{.*}}spir_func void @{{.*}}external_10{{.*}}() #{{[0-9]+}}{{.*}} !intel_reqd_sub_group_size ![[TEN:[0-9]+]]

SYCL_EXTERNAL void external_default_behavior() {}
// NONE-DAG: define {{.*}}spir_func void @{{.*}}external_default_behavior{{.*}}() #{{[0-9]+}} !srcloc !{{[0-9]+}} !sycl_fixed_targets !{{[0-9]+}} {
// NONE-DAG: define {{.*}}spir_func void @{{.*}}external_default_behavior{{.*}}() #{{[0-9]+}} !sycl_fixed_targets !{{[0-9]+}} {
// PRIM_DEF-DAG: define {{.*}}spir_func void @{{.*}}external_default_behavior{{.*}}() #{{[0-9]+}}{{.*}} !intel_reqd_sub_group_size ![[PRIMARY]]
// TEN_DEF-DAG: define {{.*}}spir_func void @{{.*}}external_default_behavior{{.*}}() #{{[0-9]+}}{{.*}} !intel_reqd_sub_group_size ![[TEN]]

void default_behavior() {
kernel_single_task<class Kernel1>([]() {
});
}
// NONE-DAG: define {{.*}}spir_kernel void @{{.*}}Kernel1() #{{[0-9]+}} !srcloc !{{[0-9]+}} !kernel_arg_buffer_location !{{[0-9]+}} !sycl_fixed_targets !{{[0-9]+}} {
// NONE-DAG: define {{.*}}spir_kernel void @{{.*}}Kernel1() #{{[0-9]+}} !kernel_arg_buffer_location !{{[0-9]+}} !sycl_fixed_targets !{{[0-9]+}} {
// PRIM_DEF-DAG: define {{.*}}spir_kernel void @{{.*}}Kernel1() #{{[0-9]+}}{{.*}} !intel_reqd_sub_group_size ![[PRIMARY]]
// TEN_DEF-DAG: define {{.*}}spir_kernel void @{{.*}}Kernel1() #{{[0-9]+}}{{.*}} !intel_reqd_sub_group_size ![[TEN]]

Expand Down
11 changes: 8 additions & 3 deletions llvm/lib/SYCLLowerIR/CleanupSYCLMetadata.cpp
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 Down Expand Up @@ -65,7 +64,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);

// Cleanup no longer needed function metadata.
for (auto &F : M) {
if (F.getMetadata("srcloc"))
F.setMetadata("srcloc", nullptr);
}

return PreservedAnalyses::all();
}
Expand Down
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 srcloc metadata
; 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
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
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
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]+}}}
Loading