Skip to content

Commit ddb561e

Browse files
aelovikov-intelaelovikov-intel
authored
Merge branch 'sycl' into handler_impl
2 parents fd4300f + 6e55f3b commit ddb561e

34 files changed

+1400
-1403
lines changed

‎llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td‎

Lines changed: 45 additions & 34 deletions
Original file line numberDiff line numberDiff line change
@@ -321,38 +321,49 @@ defvar HipSubgroupSizesGCN5 = [64]; // gfx900-gfx906 GCN5.0 (known as "Vega"
321321
defvar HipSubgroupSizesRDNA = [32]; // gfxX10-gfx11 (encapsulates RDNA1..3), (wave64 mode available but not used).
322322
defvar HipSubgroupSizesCDNA = [64]; // gfx908, gfx90a (encapsulates CDNA1..2)
323323

324-
defvar HipMinAspects = [AspectGpu, AspectFp64, AspectOnline_compiler, AspectOnline_linker, AspectQueue_profiling,
325-
AspectExt_intel_pci_address, AspectExt_intel_max_mem_bandwidth, AspectExt_intel_device_id,
326-
AspectExt_intel_memory_clock_rate, AspectExt_intel_memory_bus_width, AspectExt_intel_free_memory];
324+
defvar HipMinAspects = [AspectGpu, AspectFp16, AspectFp64,
325+
AspectOnline_compiler, AspectOnline_linker, AspectQueue_profiling,
326+
AspectExt_intel_pci_address, AspectExt_intel_max_mem_bandwidth,
327+
AspectExt_intel_device_id, AspectExt_intel_memory_clock_rate,
328+
AspectExt_intel_memory_bus_width, AspectExt_intel_free_memory];
327329

330+
defvar HipUSMAspects = !listremove(AllUSMAspects, [AspectUsm_system_allocations]);
331+
defvar HipGraphAspects = [AspectExt_oneapi_graph, AspectExt_oneapi_limited_graph];
328332
// The following AMDGCN targets are ordered based on their ROCm driver support:
329333
//
330334
// Officially supported:
331-
def : HipTargetInfo<"amd_gpu_gfx908", !listconcat(HipMinAspects, AllUSMAspects,
332-
[AspectExt_intel_device_info_uuid, AspectExt_oneapi_graph, AspectExt_oneapi_limited_graph]), HipSubgroupSizesCDNA>;
333-
def : HipTargetInfo<"amd_gpu_gfx90a", !listconcat(HipMinAspects, AllUSMAspects,
334-
[AspectExt_intel_device_info_uuid, AspectExt_oneapi_graph, AspectExt_oneapi_limited_graph, AspectExt_oneapi_native_assert]),
335+
def : HipTargetInfo<"amd_gpu_gfx908", !listconcat(
336+
HipMinAspects, HipUSMAspects, HipGraphAspects,
337+
[AspectExt_intel_device_info_uuid]), HipSubgroupSizesCDNA>;
338+
def : HipTargetInfo<"amd_gpu_gfx90a", !listconcat(
339+
HipMinAspects, HipUSMAspects, HipGraphAspects,
340+
[AspectAtomic64, AspectExt_intel_device_info_uuid, AspectExt_oneapi_native_assert]),
335341
HipSubgroupSizesCDNA>;
336342
// TODO: Need to verify whether device-side asserts (oneapi_native_assert) are
337343
// now working for the new CDNA3 gfx940, gfx941, gfx942 GPUs and fixed for the
338344
// other supported, gfx1030 and gfx1100, RDNA3 GPUs.
339-
def : HipTargetInfo<"amd_gpu_gfx940", !listconcat(HipMinAspects, AllUSMAspects,
340-
[AspectExt_intel_device_info_uuid, AspectExt_oneapi_graph, AspectExt_oneapi_limited_graph]),
345+
def : HipTargetInfo<"amd_gpu_gfx940", !listconcat(
346+
HipMinAspects, HipUSMAspects, HipGraphAspects,
347+
[AspectExt_intel_device_info_uuid]),
341348
HipSubgroupSizesCDNA>;
342-
def : HipTargetInfo<"amd_gpu_gfx941", !listconcat(HipMinAspects, AllUSMAspects,
343-
[AspectExt_intel_device_info_uuid, AspectExt_oneapi_graph, AspectExt_oneapi_limited_graph]),
349+
def : HipTargetInfo<"amd_gpu_gfx941", !listconcat(
350+
HipMinAspects, HipUSMAspects, HipGraphAspects,
351+
[AspectExt_intel_device_info_uuid]),
344352
HipSubgroupSizesCDNA>;
345-
def : HipTargetInfo<"amd_gpu_gfx942", !listconcat(HipMinAspects, AllUSMAspects,
346-
[AspectExt_intel_device_info_uuid, AspectExt_oneapi_graph, AspectExt_oneapi_limited_graph]),
353+
def : HipTargetInfo<"amd_gpu_gfx942", !listconcat(
354+
HipMinAspects, HipUSMAspects, HipGraphAspects,
355+
[AspectExt_intel_device_info_uuid]),
347356
HipSubgroupSizesCDNA>;
348-
def : HipTargetInfo<"amd_gpu_gfx1030", !listconcat(HipMinAspects, AllUSMAspects,
349-
[AspectExt_intel_device_info_uuid, AspectExt_oneapi_graph, AspectExt_oneapi_limited_graph]),
357+
def : HipTargetInfo<"amd_gpu_gfx1030", !listconcat(
358+
HipMinAspects, HipUSMAspects, HipGraphAspects,
359+
[AspectAtomic64, AspectExt_intel_device_info_uuid]),
350360
HipSubgroupSizesRDNA>;
351-
def : HipTargetInfo<"amd_gpu_gfx1100", !listconcat(HipMinAspects, AllUSMAspects,
352-
[AspectExt_intel_device_info_uuid, AspectExt_oneapi_graph, AspectExt_oneapi_limited_graph]),
361+
def : HipTargetInfo<"amd_gpu_gfx1100", !listconcat(
362+
HipMinAspects, HipUSMAspects, HipGraphAspects,
363+
[AspectExt_intel_device_info_uuid]),
353364
HipSubgroupSizesRDNA>;
354365
// Deprecated support:
355-
def : HipTargetInfo<"amd_gpu_gfx906", !listconcat(HipMinAspects, AllUSMAspects), HipSubgroupSizesGCN5>;
366+
def : HipTargetInfo<"amd_gpu_gfx906", !listconcat(HipMinAspects, HipUSMAspects), HipSubgroupSizesGCN5>;
356367
// Unsupported (or unofficially supported):
357368
def : HipTargetInfo<"amd_gpu_gfx700", HipMinAspects, HipSubgroupSizesGCN2>;
358369
def : HipTargetInfo<"amd_gpu_gfx701", HipMinAspects, HipSubgroupSizesGCN2>;
@@ -369,23 +380,23 @@ def : HipTargetInfo<"amd_gpu_gfx900", HipMinAspects, HipSubgroupSizesGCN5>;
369380
def : HipTargetInfo<"amd_gpu_gfx902", HipMinAspects, HipSubgroupSizesGCN5>;
370381
def : HipTargetInfo<"amd_gpu_gfx904", HipMinAspects, HipSubgroupSizesGCN5>;
371382
def : HipTargetInfo<"amd_gpu_gfx909", HipMinAspects, HipSubgroupSizesGCN5>;
372-
def : HipTargetInfo<"amd_gpu_gfx90c", !listconcat(HipMinAspects, AllUSMAspects), HipSubgroupSizesGCN5>;
373-
def : HipTargetInfo<"amd_gpu_gfx1010", !listconcat(HipMinAspects, AllUSMAspects), HipSubgroupSizesRDNA>;
374-
def : HipTargetInfo<"amd_gpu_gfx1011", !listconcat(HipMinAspects, AllUSMAspects), HipSubgroupSizesRDNA>;
375-
def : HipTargetInfo<"amd_gpu_gfx1012", !listconcat(HipMinAspects, AllUSMAspects), HipSubgroupSizesRDNA>;
376-
def : HipTargetInfo<"amd_gpu_gfx1013", !listconcat(HipMinAspects, AllUSMAspects), HipSubgroupSizesRDNA>;
377-
def : HipTargetInfo<"amd_gpu_gfx1031", !listconcat(!listremove(HipMinAspects, [AspectExt_intel_free_memory]), AllUSMAspects),
383+
def : HipTargetInfo<"amd_gpu_gfx90c", !listconcat(HipMinAspects, HipUSMAspects), HipSubgroupSizesGCN5>;
384+
def : HipTargetInfo<"amd_gpu_gfx1010", !listconcat(HipMinAspects, HipUSMAspects), HipSubgroupSizesRDNA>;
385+
def : HipTargetInfo<"amd_gpu_gfx1011", !listconcat(HipMinAspects, HipUSMAspects), HipSubgroupSizesRDNA>;
386+
def : HipTargetInfo<"amd_gpu_gfx1012", !listconcat(HipMinAspects, HipUSMAspects), HipSubgroupSizesRDNA>;
387+
def : HipTargetInfo<"amd_gpu_gfx1013", !listconcat(HipMinAspects, HipUSMAspects), HipSubgroupSizesRDNA>;
388+
def : HipTargetInfo<"amd_gpu_gfx1031", !listconcat(!listremove(HipMinAspects, [AspectExt_intel_free_memory]), HipUSMAspects),
378389
HipSubgroupSizesRDNA>;
379-
def : HipTargetInfo<"amd_gpu_gfx1032", !listconcat(HipMinAspects, AllUSMAspects), HipSubgroupSizesRDNA>;
380-
def : HipTargetInfo<"amd_gpu_gfx1033", !listconcat(HipMinAspects, AllUSMAspects), HipSubgroupSizesRDNA>;
381-
def : HipTargetInfo<"amd_gpu_gfx1034", !listconcat(HipMinAspects, AllUSMAspects), HipSubgroupSizesRDNA>;
382-
def : HipTargetInfo<"amd_gpu_gfx1035", !listconcat(HipMinAspects, AllUSMAspects), HipSubgroupSizesRDNA>;
383-
def : HipTargetInfo<"amd_gpu_gfx1036", !listconcat(HipMinAspects, AllUSMAspects), HipSubgroupSizesRDNA>;
384-
def : HipTargetInfo<"amd_gpu_gfx1101", !listconcat(HipMinAspects, AllUSMAspects), HipSubgroupSizesRDNA>;
385-
def : HipTargetInfo<"amd_gpu_gfx1102", !listconcat(HipMinAspects, AllUSMAspects), HipSubgroupSizesRDNA>;
386-
def : HipTargetInfo<"amd_gpu_gfx1103", !listconcat(HipMinAspects, AllUSMAspects), HipSubgroupSizesRDNA>;
387-
def : HipTargetInfo<"amd_gpu_gfx1150", !listconcat(HipMinAspects, AllUSMAspects), HipSubgroupSizesRDNA>;
388-
def : HipTargetInfo<"amd_gpu_gfx1151", !listconcat(HipMinAspects, AllUSMAspects), HipSubgroupSizesRDNA>;
390+
def : HipTargetInfo<"amd_gpu_gfx1032", !listconcat(HipMinAspects, HipUSMAspects), HipSubgroupSizesRDNA>;
391+
def : HipTargetInfo<"amd_gpu_gfx1033", !listconcat(HipMinAspects, HipUSMAspects), HipSubgroupSizesRDNA>;
392+
def : HipTargetInfo<"amd_gpu_gfx1034", !listconcat(HipMinAspects, HipUSMAspects), HipSubgroupSizesRDNA>;
393+
def : HipTargetInfo<"amd_gpu_gfx1035", !listconcat(HipMinAspects, HipUSMAspects), HipSubgroupSizesRDNA>;
394+
def : HipTargetInfo<"amd_gpu_gfx1036", !listconcat(HipMinAspects, HipUSMAspects), HipSubgroupSizesRDNA>;
395+
def : HipTargetInfo<"amd_gpu_gfx1101", !listconcat(HipMinAspects, HipUSMAspects), HipSubgroupSizesRDNA>;
396+
def : HipTargetInfo<"amd_gpu_gfx1102", !listconcat(HipMinAspects, HipUSMAspects), HipSubgroupSizesRDNA>;
397+
def : HipTargetInfo<"amd_gpu_gfx1103", !listconcat(HipMinAspects, HipUSMAspects), HipSubgroupSizesRDNA>;
398+
def : HipTargetInfo<"amd_gpu_gfx1150", !listconcat(HipMinAspects, HipUSMAspects), HipSubgroupSizesRDNA>;
399+
def : HipTargetInfo<"amd_gpu_gfx1151", !listconcat(HipMinAspects, HipUSMAspects), HipSubgroupSizesRDNA>;
389400
// TBA
390401
def : HipTargetInfo<"amd_gpu_gfx1200", [], []>; // RDNA 4
391402
def : HipTargetInfo<"amd_gpu_gfx1201", [], []>; // RDNA 4

‎llvm/include/llvm/Transforms/Instrumentation/SPIRVSanitizerCommonUtils.h‎

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -12,6 +12,7 @@
1212
#ifndef LLVM_TRANSFORMS_INSTRUMENTATION_SPIRVSANITIZERCOMMONUTILS_H
1313
#define LLVM_TRANSFORMS_INSTRUMENTATION_SPIRVSANITIZERCOMMONUTILS_H
1414

15+
#include "llvm/ADT/SmallString.h"
1516
#include "llvm/IR/Constants.h"
1617
#include "llvm/IR/DerivedTypes.h"
1718
#include "llvm/IR/Type.h"
@@ -36,6 +37,11 @@ bool isJointMatrixAccess(Value *V);
3637
// it has been used.
3738
void getFunctionsOfUser(User *User, SmallVectorImpl<Function *> &Functions);
3839

40+
// Compute MD5 hash for kernel metadata global as unique id.
41+
SmallString<128>
42+
computeKernelMetadataUniqueId(StringRef Prefix,
43+
SmallVectorImpl<uint8_t> &KernelNamesBytes);
44+
3945
} // namespace llvm
4046

4147
#endif // LLVM_TRANSFORMS_INSTRUMENTATION_SPIRVSANITIZERCOMMONUTILS_H

‎llvm/lib/SYCLLowerIR/LowerWGLocalMemory.cpp‎

Lines changed: 5 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -14,6 +14,7 @@
1414
#include "llvm/IR/IRBuilder.h"
1515
#include "llvm/IR/InstIterator.h"
1616
#include "llvm/Pass.h"
17+
#include "llvm/SYCLLowerIR/SYCLUtils.h"
1718
#include "llvm/TargetParser/Triple.h"
1819
#include "llvm/Transforms/Utils/Cloning.h"
1920

@@ -110,7 +111,11 @@ static bool inlineGroupLocalMemoryFunc(Module &M) {
110111
for (auto *U : make_early_inc_range(F->users())) {
111112
auto *CI = cast<CallInst>(U);
112113
auto *Caller = CI->getFunction();
114+
// Frontend propagates sycl-forceinline attribute to SYCL_EXTERNAL
115+
// function which directly calls group_local_memory_for_overwrite.
116+
// Don't inline the SYCL_EXTERNAL function.
113117
if (Caller->hasFnAttribute("sycl-forceinline") &&
118+
!sycl::utils::isSYCLExternalFunction(Caller) &&
114119
Visited.insert(Caller).second)
115120
WorkList.push_back(Caller);
116121
if (F != ALMFunc) {
@@ -119,8 +124,6 @@ static bool inlineGroupLocalMemoryFunc(Module &M) {
119124
assert(Result.isSuccess() && "inlining failed");
120125
}
121126
}
122-
if (F != ALMFunc)
123-
F->eraseFromParent();
124127
}
125128

126129
return !Visited.empty();

‎llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp‎

Lines changed: 5 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1402,6 +1402,7 @@ static void ExtendSpirKernelArgs(Module &M, FunctionAnalysisManager &FAM,
14021402
bool HasESIMD) {
14031403
SmallVector<Function *> SpirFixupKernels;
14041404
SmallVector<Constant *, 8> SpirKernelsMetadata;
1405+
SmallVector<uint8_t, 256> KernelNamesBytes;
14051406

14061407
const auto &DL = M.getDataLayout();
14071408
Type *IntptrTy = DL.getIntPtrType(M.getContext());
@@ -1438,6 +1439,7 @@ static void ExtendSpirKernelArgs(Module &M, FunctionAnalysisManager &FAM,
14381439
SpirFixupKernels.emplace_back(&F);
14391440

14401441
auto KernelName = F.getName();
1442+
KernelNamesBytes.append(KernelName.begin(), KernelName.end());
14411443
auto *KernelNameGV = GetOrCreateGlobalString(
14421444
M, "__asan_kernel", KernelName, kSpirOffloadConstantAS);
14431445
SpirKernelsMetadata.emplace_back(ConstantStruct::get(
@@ -1459,8 +1461,9 @@ static void ExtendSpirKernelArgs(Module &M, FunctionAnalysisManager &FAM,
14591461
"sycl-device-global-size", std::to_string(DL.getTypeAllocSize(ArrayTy)));
14601462
AsanSpirKernelMetadata->addAttribute("sycl-device-image-scope");
14611463
AsanSpirKernelMetadata->addAttribute("sycl-host-access", "0"); // read only
1462-
AsanSpirKernelMetadata->addAttribute("sycl-unique-id",
1463-
"_Z20__AsanKernelMetadata");
1464+
AsanSpirKernelMetadata->addAttribute(
1465+
"sycl-unique-id",
1466+
computeKernelMetadataUniqueId("__AsanKernelMetadata", KernelNamesBytes));
14641467
AsanSpirKernelMetadata->setDSOLocal(true);
14651468

14661469
// Handle SpirFixupKernels

‎llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp‎

Lines changed: 5 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1171,6 +1171,7 @@ void MemorySanitizerOnSpirv::instrumentPrivateArguments(
11711171
// kernel
11721172
void MemorySanitizerOnSpirv::instrumentKernelsMetadata() {
11731173
SmallVector<Constant *, 8> SpirKernelsMetadata;
1174+
SmallVector<uint8_t, 256> KernelNamesBytes;
11741175

11751176
// SpirKernelsMetadata only saves fixed kernels, and is described by
11761177
// following structure:
@@ -1189,6 +1190,7 @@ void MemorySanitizerOnSpirv::instrumentKernelsMetadata() {
11891190
continue;
11901191

11911192
auto KernelName = F.getName();
1193+
KernelNamesBytes.append(KernelName.begin(), KernelName.end());
11921194
auto *KernelNameGV = getOrCreateGlobalString("__msan_kernel", KernelName,
11931195
kSpirOffloadConstantAS);
11941196
SpirKernelsMetadata.emplace_back(ConstantStruct::get(
@@ -1213,8 +1215,9 @@ void MemorySanitizerOnSpirv::instrumentKernelsMetadata() {
12131215
MsanSpirKernelMetadata->addAttribute("sycl-device-image-scope");
12141216
MsanSpirKernelMetadata->addAttribute("sycl-host-access",
12151217
"0"); // read only
1216-
MsanSpirKernelMetadata->addAttribute("sycl-unique-id",
1217-
"_Z20__MsanKernelMetadata");
1218+
MsanSpirKernelMetadata->addAttribute(
1219+
"sycl-unique-id",
1220+
computeKernelMetadataUniqueId("__MsanKernelMetadata", KernelNamesBytes));
12181221
MsanSpirKernelMetadata->setDSOLocal(true);
12191222
}
12201223

‎llvm/lib/Transforms/Instrumentation/SPIRVSanitizerCommonUtils.cpp‎

Lines changed: 13 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -12,6 +12,7 @@
1212

1313
#include "llvm/Transforms/Instrumentation/SPIRVSanitizerCommonUtils.h"
1414
#include "llvm/IR/Instructions.h"
15+
#include "llvm/Support/MD5.h"
1516

1617
using namespace llvm;
1718

@@ -68,4 +69,16 @@ void getFunctionsOfUser(User *User, SmallVectorImpl<Function *> &Functions) {
6869
}
6970
}
7071

72+
SmallString<128>
73+
computeKernelMetadataUniqueId(StringRef Prefix,
74+
SmallVectorImpl<uint8_t> &KernelNamesBytes) {
75+
MD5 Hash;
76+
SmallString<32> UniqueIdSuffix;
77+
SmallString<128> UniqueId(Prefix);
78+
auto R = Hash.hash(KernelNamesBytes);
79+
Hash.stringifyResult(R, UniqueIdSuffix);
80+
UniqueId.append(UniqueIdSuffix);
81+
return UniqueId;
82+
}
83+
7184
} // namespace llvm

‎llvm/lib/Transforms/Instrumentation/ThreadSanitizer.cpp‎

Lines changed: 6 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -37,6 +37,7 @@
3737
#include "llvm/IR/LLVMContext.h"
3838
#include "llvm/IR/Metadata.h"
3939
#include "llvm/IR/Module.h"
40+
#include "llvm/IR/ModuleSummaryIndex.h"
4041
#include "llvm/IR/Type.h"
4142
#include "llvm/ProfileData/InstrProf.h"
4243
#include "llvm/Support/CommandLine.h"
@@ -674,6 +675,7 @@ void ThreadSanitizerOnSpirv::initializeKernelCallerMap(Function *F) {
674675

675676
void ThreadSanitizerOnSpirv::instrumentKernelsMetadata() {
676677
SmallVector<Constant *, 8> SpirKernelsMetadata;
678+
SmallVector<uint8_t, 256> KernelNamesBytes;
677679

678680
// SpirKernelsMetadata only saves fixed kernels, and is described by
679681
// following structure:
@@ -687,6 +689,7 @@ void ThreadSanitizerOnSpirv::instrumentKernelsMetadata() {
687689

688690
if (isSupportedSPIRKernel(F)) {
689691
auto KernelName = F.getName();
692+
KernelNamesBytes.append(KernelName.begin(), KernelName.end());
690693
auto *KernelNameGV = GetOrCreateGlobalString("__tsan_kernel", KernelName,
691694
kSpirOffloadConstantAS);
692695
SpirKernelsMetadata.emplace_back(ConstantStruct::get(
@@ -709,8 +712,9 @@ void ThreadSanitizerOnSpirv::instrumentKernelsMetadata() {
709712
"sycl-device-global-size", std::to_string(DL.getTypeAllocSize(ArrayTy)));
710713
TsanSpirKernelMetadata->addAttribute("sycl-device-image-scope");
711714
TsanSpirKernelMetadata->addAttribute("sycl-host-access", "0"); // read only
712-
TsanSpirKernelMetadata->addAttribute("sycl-unique-id",
713-
"_Z20__TsanKernelMetadata");
715+
TsanSpirKernelMetadata->addAttribute(
716+
"sycl-unique-id",
717+
computeKernelMetadataUniqueId("__TsanKernelMetadata", KernelNamesBytes));
714718
TsanSpirKernelMetadata->setDSOLocal(true);
715719
}
716720

‎llvm/test/Instrumentation/AddressSanitizer/SPIRV/extend_launch_info_arg.ll‎

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -4,6 +4,7 @@ target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:
44
target triple = "spir64-unknown-unknown"
55

66
; CHECK: @__AsanKernelMetadata = appending dso_local local_unnamed_addr addrspace(1) global
7+
; CHECK-SAME: [[ATTR0:#[0-9]+]]
78
; CHECK: @__AsanLaunchInfo = external addrspace(3) global ptr addrspace(1)
89

910
define spir_kernel void @sycl_kernel1() #0 {
@@ -25,4 +26,4 @@ entry:
2526
attributes #0 = { sanitize_address }
2627
;; sycl-device-global-size = 16 * 2
2728
;; sycl-host-access = 0 read-only
28-
; CHECK: attributes #{{.*}} = { "sycl-device-global-size"="32" "sycl-device-image-scope" "sycl-host-access"="0" "sycl-unique-id"="_Z20__AsanKernelMetadata" }
29+
; CHECK: attributes [[ATTR0]] = { "sycl-device-global-size"="32" "sycl-device-image-scope" "sycl-host-access"="0" "sycl-unique-id"="__AsanKernelMetadata833c47834a0b74946e370c23c39607cc" }

‎llvm/test/Instrumentation/MemorySanitizer/SPIRV/instrument_global_address_space.ll‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -31,4 +31,4 @@ entry:
3131
}
3232

3333
; CHECK: attributes [[ATTR0]]
34-
; CHECK-SAME: "sycl-device-global-size"="32" "sycl-device-image-scope" "sycl-host-access"="0" "sycl-unique-id"="_Z20__MsanKernelMetadata"
34+
; CHECK-SAME: "sycl-device-global-size"="32" "sycl-device-image-scope" "sycl-host-access"="0" "sycl-unique-id"="__MsanKernelMetadata3ff767e9a7a43f1f3968062dbb4ee3b4"

‎llvm/test/Instrumentation/ThreadSanitizer/SPIRV/kernel_metadata.ll‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -4,11 +4,13 @@ target triple = "spir64-unknown-unknown"
44

55
; CHECK-LABEL: @__TsanKernelMetadata = appending dso_local local_unnamed_addr addrspace(1) global
66
; CHECK-SAME: i64 ptrtoint (ptr addrspace(2) @__tsan_kernel to i64
7+
; CHECK-SAME: [[ATTR0:#[0-9]+]]
78

89
; Function Attrs: sanitize_thread
910
define spir_kernel void @test() #0 {
1011
entry:
1112
ret void
1213
}
1314

15+
; CHECK: attributes [[ATTR0]] = {{.*}} "sycl-unique-id"="__TsanKernelMetadata098f6bcd4621d373cade4e832627b4f6"
1416
attributes #0 = { sanitize_thread }

0 commit comments

Comments
 (0)