Skip to content

Commit 4805490

Browse files
MrSidimsMrSidims
authored
[SYCL] Add pass that remove compiler internal metadata (#11490)
Right now the pass removes only sycl_aspects and sycl_types_that_use_aspects module-level metadata, so during llvm-link these metadata nodes won't be duplicated significantly increasing module size. Other metadata like sycl_used_aspects (applied to a function) were also considered to be added to the pass, but at this moment such metadata causes no harm, so it was left untouched. --------- Signed-off-by: Sidorov, Dmitry <[email protected]>
1 parent 0f52ee6 commit 4805490

File tree

13 files changed

+178
-19
lines changed

13 files changed

+178
-19
lines changed

‎clang/lib/CodeGen/BackendUtil.cpp

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -43,6 +43,7 @@
4343
#include "llvm/Passes/PassBuilder.h"
4444
#include "llvm/Passes/PassPlugin.h"
4545
#include "llvm/Passes/StandardInstrumentations.h"
46+
#include "llvm/SYCLLowerIR/CleanupSYCLMetadata.h"
4647
#include "llvm/SYCLLowerIR/CompileTimePropertiesPass.h"
4748
#include "llvm/SYCLLowerIR/ESIMD/ESIMDVerifier.h"
4849
#include "llvm/SYCLLowerIR/ESIMD/LowerESIMD.h"
@@ -1127,6 +1128,10 @@ void EmitAssemblyHelper::RunOptimizationPipeline(
11271128
if (LangOpts.SYCLIsNativeCPU) {
11281129
MPM.addPass(PrepareSYCLNativeCPUPass());
11291130
}
1131+
1132+
// Remove SYCL metadata added by the frontend, like sycl_aspects
1133+
// Note, this pass should be at the end of the pipeline
1134+
MPM.addPass(CleanupSYCLMetadataPass());
11301135
}
11311136
}
11321137

‎llvm/include/llvm/SYCLLowerIR/CleanupSYCLMetadata.h

Lines changed: 28 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,28 @@
1+
//===---------- CleanupSYCLMetadata.h - CleanupSYCLMetadata Pass ----------===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
//
9+
// Cleanup SYCL compiler internal metadata inserted by the frontend as it will
10+
// never be used in the compilation ever again
11+
//
12+
//===----------------------------------------------------------------------===//
13+
//
14+
#ifndef LLVM_CLEANUP_SYCL_METADATA
15+
#define LLVM_CLEANUP_SYCL_METADATA
16+
17+
#include "llvm/IR/PassManager.h"
18+
19+
namespace llvm {
20+
21+
class CleanupSYCLMetadataPass : public PassInfoMixin<CleanupSYCLMetadataPass> {
22+
public:
23+
PreservedAnalyses run(Module &M, ModuleAnalysisManager &);
24+
};
25+
26+
} // namespace llvm
27+
28+
#endif // LLVM_CLEANUP_SYCL_METADATA

‎llvm/lib/Passes/PassBuilder.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -82,6 +82,7 @@
8282
#include "llvm/IR/Verifier.h"
8383
#include "llvm/IRPrinter/IRPrintingPasses.h"
8484
#include "llvm/Passes/OptimizationLevel.h"
85+
#include "llvm/SYCLLowerIR/CleanupSYCLMetadata.h"
8586
#include "llvm/SYCLLowerIR/CompileTimePropertiesPass.h"
8687
#include "llvm/SYCLLowerIR/ESIMD/ESIMDVerifier.h"
8788
#include "llvm/SYCLLowerIR/ESIMD/LowerESIMD.h"

‎llvm/lib/Passes/PassRegistry.def

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -142,6 +142,7 @@ MODULE_PASS("esimd-remove-host-code", ESIMDRemoveHostCodePass());
142142
MODULE_PASS("sycl-propagate-aspects-usage", SYCLPropagateAspectsUsagePass())
143143
MODULE_PASS("sycl-add-opt-level-attribute", SYCLAddOptLevelAttributePass())
144144
MODULE_PASS("compile-time-properties", CompileTimePropertiesPass())
145+
MODULE_PASS("cleanup-sycl-metadata", CleanupSYCLMetadataPass())
145146
#undef MODULE_PASS
146147

147148
#ifndef MODULE_PASS_WITH_PARAMS

‎llvm/lib/SYCLLowerIR/CMakeLists.txt

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -53,6 +53,7 @@ add_llvm_component_library(LLVMSYCLLowerIR
5353
ESIMD/ESIMDRemoveHostCode.cpp
5454
ESIMD/LowerESIMD.cpp
5555
ESIMD/LowerESIMDKernelAttrs.cpp
56+
CleanupSYCLMetadata.cpp
5657
CompileTimePropertiesPass.cpp
5758
DeviceGlobals.cpp
5859
ESIMD/LowerESIMDVLoadVStore.cpp

‎llvm/lib/SYCLLowerIR/CleanupSYCLMetadata.cpp

Lines changed: 45 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,45 @@
1+
//===--------- CleanupSYCLMetadata.cpp - CleanupSYCLMetadata Pass ---------===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
//
9+
// Cleanup SYCL compiler internal metadata inserted by the frontend as it will
10+
// never be used in the compilation ever again
11+
//
12+
//===----------------------------------------------------------------------===//
13+
14+
#include "llvm/SYCLLowerIR/CleanupSYCLMetadata.h"
15+
16+
#include "llvm/IR/Module.h"
17+
#include "llvm/Pass.h"
18+
19+
using namespace llvm;
20+
21+
namespace {
22+
23+
void cleanupSYCLCompilerMetadata(const Module &M, llvm::StringRef MD) {
24+
NamedMDNode *Node = M.getNamedMetadata(MD);
25+
if (!Node)
26+
return;
27+
Node->clearOperands();
28+
Node->dropAllReferences();
29+
Node->eraseFromParent();
30+
}
31+
32+
} // anonymous namespace
33+
34+
PreservedAnalyses CleanupSYCLMetadataPass::run(Module &M,
35+
ModuleAnalysisManager &MAM) {
36+
// Remove SYCL module-level metadata that will never be used again to avoid
37+
// duplication of their operands during llvm-link hence preventing
38+
// increase of the module size
39+
llvm::SmallVector<llvm::StringRef, 2> ModuleMDToRemove = {
40+
"sycl_aspects", "sycl_types_that_use_aspects"};
41+
for (const auto &MD : ModuleMDToRemove)
42+
cleanupSYCLCompilerMetadata(M, MD);
43+
44+
return PreservedAnalyses::all();
45+
}

‎llvm/test/SYCLLowerIR/CleanupSYCLCompilerInternalMetadata/aspects-md-in-two-nodes.ll

Lines changed: 23 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,23 @@
1+
; RUN: opt -passes=cleanup-sycl-metadata -S < %s | FileCheck %s
2+
;
3+
; Test checks that the pass is able to cleanup sycl_aspects and
4+
; sycl_types_that_use_aspects module metadata
5+
6+
; ModuleID = 'funny-aspects.ll'
7+
source_filename = "funny-aspects.ll"
8+
9+
%A = type { i32 }
10+
11+
define spir_kernel void @kernel() !artificial !0 {
12+
ret void
13+
}
14+
15+
; CHECK-NOT: sycl_types_that_use_aspects
16+
; CHECK-NOT: sycl_aspects
17+
; CHECK: !0 = !{!"A", i32 0}
18+
19+
!sycl_types_that_use_aspects = !{!0}
20+
!sycl_aspects = !{!1}
21+
22+
!0 = !{!"A", i32 0}
23+
!1 = !{!"fp64", i32 6}

‎llvm/test/SYCLLowerIR/CleanupSYCLCompilerInternalMetadata/multiple-aspects.ll

Lines changed: 56 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,56 @@
1+
; RUN: opt -passes=cleanup-sycl-metadata -S < %s | FileCheck %s
2+
;
3+
; Test checks that the pass is able to cleanup sycl_aspects and
4+
; sycl_types_that_use_aspects module metadata
5+
6+
; ModuleID = 'multiple-aspects.ll'
7+
source_filename = "multiple-aspects.ll"
8+
9+
%A = type { i32 }
10+
%B = type { i32 }
11+
%C = type { i32 }
12+
%D = type { i32 }
13+
14+
define spir_func void @funcA() !sycl_used_aspects !5 {
15+
%tmp = alloca %A, align 8
16+
ret void
17+
}
18+
19+
define spir_func void @funcB() !sycl_used_aspects !6 {
20+
%tmp = alloca %B, align 8
21+
call spir_func void @funcA()
22+
ret void
23+
}
24+
25+
define spir_func void @funcC() !sycl_used_aspects !7 {
26+
%tmp = alloca %C, align 8
27+
call spir_func void @funcB()
28+
ret void
29+
}
30+
31+
define spir_func void @funcD() !sycl_used_aspects !8 {
32+
%tmp = alloca %D, align 8
33+
call spir_func void @funcC()
34+
ret void
35+
}
36+
37+
define spir_kernel void @kernel() !sycl_used_aspects !8 !sycl_fixed_targets !9 {
38+
call spir_func void @funcD()
39+
ret void
40+
}
41+
42+
; CHECK-NOT: sycl_types_that_use_aspects
43+
; CHECK-NOT: sycl_aspects
44+
!sycl_types_that_use_aspects = !{!0, !1, !2, !3}
45+
!sycl_aspects = !{!4}
46+
47+
!0 = !{!"A", i32 0}
48+
!1 = !{!"B", i32 1}
49+
!2 = !{!"C", i32 2}
50+
!3 = !{!"D", i32 3, i32 4}
51+
!4 = !{!"fp64", i32 6}
52+
!5 = !{i32 0}
53+
!6 = !{i32 1, i32 0}
54+
!7 = !{i32 2, i32 1, i32 0}
55+
!8 = !{i32 0, i32 1, i32 2, i32 3, i32 4}
56+
!9 = !{}

‎sycl/test/extensions/properties/properties_kernel_device_has.cpp

Lines changed: 14 additions & 15 deletions
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// RUN: %clangxx -fsycl-device-only -S -Xclang -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK-IR
1+
// RUN: %clangxx -fsycl-device-only -S -Xclang -emit-llvm -Xclang -disable-llvm-passes %s -o - | FileCheck %s --check-prefix CHECK-IR
22
// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify %s
33
// expected-no-diagnostics
44

@@ -51,12 +51,12 @@ int main() {
5151
// CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel5(){{.*}} #[[DHAttr2]]
5252
Q.parallel_for<class WGSizeKernel5>(R1, {Ev}, Props, [](id<1>) {});
5353

54-
// CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel6{{.*}}{{.*}} #[[DHAttr3:[0-9]+]]
54+
// CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel6{{.*}}{{.*}} #[[DHAttr2:[0-9]+]]
5555
Q.parallel_for<class WGSizeKernel6>(R1, Props, Redu1, [](id<1>, auto &) {});
56-
// CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel7{{.*}}{{.*}} #[[DHAttr3]]
56+
// CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel7{{.*}}{{.*}} #[[DHAttr2]]
5757
Q.parallel_for<class WGSizeKernel7>(R1, Ev, Props, Redu1,
5858
[](id<1>, auto &) {});
59-
// CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel8{{.*}}{{.*}} #[[DHAttr3]]
59+
// CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel8{{.*}}{{.*}} #[[DHAttr2]]
6060
Q.parallel_for<class WGSizeKernel8>(R1, {Ev}, Props, Redu1,
6161
[](id<1>, auto &) {});
6262

@@ -67,23 +67,23 @@ int main() {
6767
// CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel11(){{.*}} #[[DHAttr2]]
6868
Q.parallel_for<class WGSizeKernel11>(NDR1, {Ev}, Props, [](nd_item<1>) {});
6969

70-
// CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel12{{.*}}{{.*}} #[[DHAttr3]]
70+
// CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel12{{.*}}{{.*}} #[[DHAttr2]]
7171
Q.parallel_for<class WGSizeKernel12>(NDR1, Props, Redu1,
7272
[](nd_item<1>, auto &) {});
73-
// CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel13{{.*}}{{.*}} #[[DHAttr3]]
73+
// CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel13{{.*}}{{.*}} #[[DHAttr2]]
7474
Q.parallel_for<class WGSizeKernel13>(NDR1, Ev, Props, Redu1,
7575
[](nd_item<1>, auto &) {});
76-
// CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel14{{.*}}{{.*}} #[[DHAttr3]]
76+
// CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel14{{.*}}{{.*}} #[[DHAttr2]]
7777
Q.parallel_for<class WGSizeKernel14>(NDR1, {Ev}, Props, Redu1,
7878
[](nd_item<1>, auto &) {});
7979

80-
// CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel15{{.*}}{{.*}} #[[DHAttr3]]
80+
// CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel15{{.*}}{{.*}} #[[DHAttr2]]
8181
Q.parallel_for<class WGSizeKernel15>(NDR1, Props, Redu1, Redu2,
8282
[](nd_item<1>, auto &, auto &) {});
83-
// CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel16{{.*}}{{.*}} #[[DHAttr3]]
83+
// CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel16{{.*}}{{.*}} #[[DHAttr2]]
8484
Q.parallel_for<class WGSizeKernel16>(NDR1, Ev, Props, Redu1, Redu2,
8585
[](nd_item<1>, auto &, auto &) {});
86-
// CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel17{{.*}}{{.*}} #[[DHAttr3]]
86+
// CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel17{{.*}}{{.*}} #[[DHAttr2]]
8787
Q.parallel_for<class WGSizeKernel17>(NDR1, {Ev}, Props, Redu1, Redu2,
8888
[](nd_item<1>, auto &, auto &) {});
8989

@@ -97,7 +97,7 @@ int main() {
9797
CGH.parallel_for<class WGSizeKernel19>(R1, Props, [](id<1>) {});
9898
});
9999

100-
// CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel20{{.*}}{{.*}} #[[DHAttr3]]
100+
// CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel20{{.*}}{{.*}} #[[DHAttr2]]
101101
Q.submit([&](handler &CGH) {
102102
CGH.parallel_for<class WGSizeKernel20>(R1, Props, Redu1,
103103
[](id<1>, auto &) {});
@@ -108,19 +108,19 @@ int main() {
108108
CGH.parallel_for<class WGSizeKernel21>(NDR1, Props, [](nd_item<1>) {});
109109
});
110110

111-
// CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel22{{.*}}{{.*}} #[[DHAttr3]]
111+
// CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel22{{.*}}{{.*}} #[[DHAttr2]]
112112
Q.submit([&](handler &CGH) {
113113
CGH.parallel_for<class WGSizeKernel22>(NDR1, Props, Redu1,
114114
[](nd_item<1>, auto &) {});
115115
});
116116

117-
// CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel23{{.*}}{{.*}} #[[DHAttr3]]
117+
// CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel23{{.*}}{{.*}} #[[DHAttr2]]
118118
Q.submit([&](handler &CGH) {
119119
CGH.parallel_for<class WGSizeKernel23>(NDR1, Props, Redu1, Redu2,
120120
[](nd_item<1>, auto &, auto &) {});
121121
});
122122

123-
// CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel24(){{.*}} #[[DHAttr3]]
123+
// CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel24(){{.*}} #[[DHAttr2]]
124124
Q.submit([&](handler &CGH) {
125125
CGH.parallel_for_work_group<class WGSizeKernel24>(
126126
R1, Props,
@@ -167,4 +167,3 @@ int main() {
167167

168168
// CHECK-IR-DAG: attributes #[[DHAttr1]] = { {{.*}}"sycl-device-has"="[[ext_oneapi_cuda_async_barrier_ASPECT_MD]],[[ext_oneapi_bfloat16_math_functions_ASPECT_MD]],[[custom_ASPECT_MD]],[[fp16_ASPECT_MD]],[[fp64_ASPECT_MD]],[[image_ASPECT_MD]],[[online_compiler_ASPECT_MD]],[[online_linker_ASPECT_MD]],[[queue_profiling_ASPECT_MD]],[[usm_device_allocations_ASPECT_MD]],[[usm_system_allocations_ASPECT_MD]],[[ext_intel_pci_address_ASPECT_MD]],[[cpu_ASPECT_MD]],[[gpu_ASPECT_MD]],[[accelerator_ASPECT_MD]],[[ext_intel_gpu_eu_count_ASPECT_MD]],[[ext_intel_gpu_subslices_per_slice_ASPECT_MD]],[[ext_intel_gpu_eu_count_per_subslice_ASPECT_MD]],[[ext_intel_max_mem_bandwidth_ASPECT_MD]],[[ext_intel_mem_channel_ASPECT_MD]],[[usm_atomic_host_allocations_ASPECT_MD]],[[usm_atomic_shared_allocations_ASPECT_MD]],[[atomic64_ASPECT_MD]],[[ext_intel_device_info_uuid_ASPECT_MD]],[[ext_oneapi_srgb_ASPECT_MD]],[[ext_intel_gpu_eu_simd_width_ASPECT_MD]],[[ext_intel_gpu_slices_ASPECT_MD]],[[ext_oneapi_native_assert_ASPECT_MD]],[[host_debuggable_ASPECT_MD]],[[ext_intel_gpu_hw_threads_per_eu_ASPECT_MD]],[[usm_host_allocations_ASPECT_MD]],[[usm_shared_allocations_ASPECT_MD]],[[ext_intel_free_memory_ASPECT_MD]],[[ext_intel_device_id_ASPECT_MD]]"
169169
// CHECK-IR-DAG: attributes #[[DHAttr2]] = { {{.*}}"sycl-device-has"="[[ext_oneapi_cuda_async_barrier_ASPECT_MD]],[[ext_oneapi_bfloat16_math_functions_ASPECT_MD]],[[custom_ASPECT_MD]],[[fp16_ASPECT_MD]],[[fp64_ASPECT_MD]],[[image_ASPECT_MD]],[[online_compiler_ASPECT_MD]],[[online_linker_ASPECT_MD]],[[queue_profiling_ASPECT_MD]],[[usm_device_allocations_ASPECT_MD]],[[usm_system_allocations_ASPECT_MD]],[[ext_intel_pci_address_ASPECT_MD]],[[cpu_ASPECT_MD]],[[gpu_ASPECT_MD]],[[accelerator_ASPECT_MD]],[[ext_intel_gpu_eu_count_ASPECT_MD]],[[ext_intel_gpu_subslices_per_slice_ASPECT_MD]],[[ext_intel_gpu_eu_count_per_subslice_ASPECT_MD]],[[ext_intel_max_mem_bandwidth_ASPECT_MD]],[[ext_intel_mem_channel_ASPECT_MD]],[[usm_atomic_host_allocations_ASPECT_MD]],[[usm_atomic_shared_allocations_ASPECT_MD]],[[atomic64_ASPECT_MD]],[[ext_intel_device_info_uuid_ASPECT_MD]],[[ext_oneapi_srgb_ASPECT_MD]],[[ext_intel_gpu_eu_simd_width_ASPECT_MD]],[[ext_intel_gpu_slices_ASPECT_MD]],[[ext_oneapi_native_assert_ASPECT_MD]],[[host_debuggable_ASPECT_MD]],[[ext_intel_gpu_hw_threads_per_eu_ASPECT_MD]],[[usm_host_allocations_ASPECT_MD]],[[usm_shared_allocations_ASPECT_MD]],[[ext_intel_free_memory_ASPECT_MD]],[[ext_intel_device_id_ASPECT_MD]]"
170-
// CHECK-IR-DAG: attributes #[[DHAttr3]] = { {{.*}}"sycl-device-has"="[[ext_oneapi_cuda_async_barrier_ASPECT_MD]],[[ext_oneapi_bfloat16_math_functions_ASPECT_MD]],[[custom_ASPECT_MD]],[[fp16_ASPECT_MD]],[[fp64_ASPECT_MD]],[[image_ASPECT_MD]],[[online_compiler_ASPECT_MD]],[[online_linker_ASPECT_MD]],[[queue_profiling_ASPECT_MD]],[[usm_device_allocations_ASPECT_MD]],[[usm_system_allocations_ASPECT_MD]],[[ext_intel_pci_address_ASPECT_MD]],[[cpu_ASPECT_MD]],[[gpu_ASPECT_MD]],[[accelerator_ASPECT_MD]],[[ext_intel_gpu_eu_count_ASPECT_MD]],[[ext_intel_gpu_subslices_per_slice_ASPECT_MD]],[[ext_intel_gpu_eu_count_per_subslice_ASPECT_MD]],[[ext_intel_max_mem_bandwidth_ASPECT_MD]],[[ext_intel_mem_channel_ASPECT_MD]],[[usm_atomic_host_allocations_ASPECT_MD]],[[usm_atomic_shared_allocations_ASPECT_MD]],[[atomic64_ASPECT_MD]],[[ext_intel_device_info_uuid_ASPECT_MD]],[[ext_oneapi_srgb_ASPECT_MD]],[[ext_intel_gpu_eu_simd_width_ASPECT_MD]],[[ext_intel_gpu_slices_ASPECT_MD]],[[ext_oneapi_native_assert_ASPECT_MD]],[[host_debuggable_ASPECT_MD]],[[ext_intel_gpu_hw_threads_per_eu_ASPECT_MD]],[[usm_host_allocations_ASPECT_MD]],[[usm_shared_allocations_ASPECT_MD]],[[ext_intel_free_memory_ASPECT_MD]],[[ext_intel_device_id_ASPECT_MD]]"

‎sycl/test/extensions/properties/properties_kernel_device_has_macro.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// RUN: %clangxx -fsycl-device-only -S -Xclang -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK-IR
1+
// RUN: %clangxx -fsycl-device-only -S -Xclang -emit-llvm -Xclang -disable-llvm-passes %s -o - | FileCheck %s --check-prefix CHECK-IR
22
// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify %s
33
// expected-no-diagnostics
44

‎sycl/test/optional_kernel_features/atomic_ref-atomic64-aspect.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// RUN: %clangxx %s -S -o %t.ll -fsycl-device-only
1+
// RUN: %clangxx %s -S -o %t.ll -fsycl-device-only -Xclang -disable-llvm-passes
22
// RUN: FileCheck %s --input-file %t.ll
33

44
// CHECK: !sycl_types_that_use_aspects = !{![[#MDNUM1:]], ![[#MDNUM2:]], ![[#MDNUM3:]]}

‎sycl/test/optional_kernel_features/esimd.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// RUN: %clangxx %s -S -o %t.ll -fsycl-device-only
1+
// RUN: %clangxx %s -S -o %t.ll -fsycl-device-only -Xclang -disable-llvm-passes
22
// RUN: FileCheck %s --input-file %t.ll
33

44
// CHECK: !sycl_types_that_use_aspects = !{![[#MDNUM1:]], ![[#MDNUM2:]]}

‎sycl/test/optional_kernel_features/half-aspect.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// RUN: %clangxx %s -S -o %t.ll -fsycl-device-only
1+
// RUN: %clangxx %s -S -o %t.ll -fsycl-device-only -Xclang -disable-llvm-passes
22
// RUN: FileCheck %s --input-file %t.ll
33

44
// CHECK: !sycl_types_that_use_aspects = !{![[#MDNUM:]]}

0 commit comments

Comments
 (0)