Skip to content

Commit 9f7e2ea

Browse files
HabKaffeeHabKaffee
authored
Merge branch 'sycl' into Realisation_of_sycl_link_build_compile
2 parents c70286a + 1df0038 commit 9f7e2ea

File tree

168 files changed

+2004
-962
lines changed

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

168 files changed

+2004
-962
lines changed

‎.github/workflows/sycl_nightly.yml

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -70,5 +70,5 @@ jobs:
7070
build_cache_root: "/__w/"
7171
build_cache_suffix: new_pm
7272
build_artifact_suffix: new_pm
73-
build_configure_extra_args: '--hip --hip-amd-arch=gfx906 --cuda --cmake-opt=-DLLVM_ENABLE_NEW_PASS_MANAGER=ON'
73+
build_configure_extra_args: '--hip --cuda --cmake-opt=-DLLVM_ENABLE_NEW_PASS_MANAGER=ON'
7474
lts_config: "hip_amdgpu;ocl_x64"

‎.github/workflows/sycl_stale_issues.yml

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -17,3 +17,5 @@ jobs:
1717
exempt-issue-labels: 'confirmed,hip,cuda,enhancement,help wanted,upstream'
1818
stale-issue-label: 'stale'
1919
exempt-all-issue-assignees: true
20+
operations-per-run: 200
21+

‎.github/workflows/sycl_windows_build_and_test.yml

Lines changed: 4 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -17,6 +17,7 @@ jobs:
1717
run: |
1818
choco install -y cuda --version 11.6.0.51123
1919
choco install -y ninja
20+
choco install -y sccache --version 0.2.15
2021
refreshenv
2122
echo CUDA_PATH=%CUDA_PATH%
2223
echo CUDA_PATH=%CUDA_PATH% >> %GITHUB_ENV%
@@ -49,6 +50,8 @@ jobs:
4950
--cmake-opt="-DCMAKE_C_COMPILER=cl" ^
5051
--cmake-opt="-DCMAKE_CXX_COMPILER=cl" ^
5152
--cmake-opt="-DCMAKE_INSTALL_PREFIX=%GITHUB_WORKSPACE%\install" ^
53+
--cmake-opt="-DCMAKE_CXX_COMPILER_LAUNCHER=sccache" ^
54+
--cmake-opt="-DCMAKE_C_COMPILER_LAUNCHER=sccache" ^
5255
--cuda
5356
- name: Build
5457
shell: cmd
@@ -64,4 +67,4 @@ jobs:
6467
uses: actions/upload-artifact@v2
6568
with:
6669
name: sycl_windows_default
67-
path: install/**/*
70+
path: install/**/*

‎clang/include/clang/Basic/Attr.td

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1347,7 +1347,7 @@ def SYCLScope : Attr {
13471347
def SYCLDeviceIndirectlyCallable : InheritableAttr {
13481348
let Spellings = [ CXX11<"intel", "device_indirectly_callable"> ];
13491349
let Subjects = SubjectList<[Function]>;
1350-
let LangOpts = [SYCLIsDevice];
1350+
let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost];
13511351
let Documentation = [SYCLDeviceIndirectlyCallableDocs];
13521352
}
13531353

‎clang/test/SemaSYCL/device-indirectly-callable-attr.cpp

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,8 +1,9 @@
11
// RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -verify %s
22
// RUN: not %clang_cc1 -fsycl-is-device -ast-dump %s | FileCheck %s
33
// RUN: %clang_cc1 -verify -DNO_SYCL %s
4+
// RUN: %clang_cc1 -fsycl-is-host -fsyntax-only -verify -DSYCL_HOST %s
45

5-
#ifndef NO_SYCL
6+
#if !defined(NO_SYCL) || defined(SYCL_HOST)
67

78
[[intel::device_indirectly_callable]] // expected-warning {{'device_indirectly_callable' attribute only applies to functions}}
89
int N;

‎llvm/lib/SYCLLowerIR/CMakeLists.txt

Lines changed: 6 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -13,11 +13,12 @@ endif()
1313
if (NOT TARGET LLVMGenXIntrinsics)
1414
if (NOT DEFINED LLVMGenXIntrinsics_SOURCE_DIR)
1515
set(LLVMGenXIntrinsics_GIT_REPO https://github.com/intel/vc-intrinsics.git)
16-
# commit a9bb6d8040c43404c5fbe3694e59c503d179d19a
17-
# Author: Nikita Rudenko <[email protected]>
18-
# Date: Tue Feb 1 14:57:43 2022 +0000
19-
# Fix attributes are not forwarded for call inst with SEV
20-
set(LLVMGenXIntrinsics_GIT_TAG a9bb6d8040c43404c5fbe3694e59c503d179d19a)
16+
# commit 8b6e209fe1269a2c6470b36dfbaa0e051d2a100f (master)
17+
# Author: Konstantin Vladimirov <[email protected]>
18+
# Date: Tue Feb 8 10:47:03 2022 +0000
19+
# introducing named barrier support in adaptor pass
20+
# named barrier required for DPC++ and other customers
21+
set(LLVMGenXIntrinsics_GIT_TAG 8b6e209fe1269a2c6470b36dfbaa0e051d2a100f)
2122

2223
message(STATUS "vc-intrinsics repo is missing. Will try to download it from ${LLVMGenXIntrinsics_GIT_REPO}")
2324
include(FetchContent)

‎llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp

Lines changed: 42 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -23,6 +23,7 @@
2323
#include "llvm/Demangle/Demangle.h"
2424
#include "llvm/Demangle/ItaniumDemangle.h"
2525
#include "llvm/GenXIntrinsics/GenXIntrinsics.h"
26+
#include "llvm/GenXIntrinsics/GenXMetadata.h"
2627
#include "llvm/IR/IRBuilder.h"
2728
#include "llvm/IR/InstIterator.h"
2829
#include "llvm/IR/Instructions.h"
@@ -445,6 +446,9 @@ class ESIMDIntrinDescTable {
445446
{"raw_send2_noresult",
446447
{"raw.send2.noresult",
447448
{a(0), a(1), ai1(2), a(3), a(4), a(5), a(6), a(7)}}},
449+
{"nbarrier", {"nbarrier", {a(0), a(1), a(2)}}},
450+
{"raw_send_nbarrier_signal",
451+
{"raw.send.noresult", {a(0), ai1(4), a(1), a(2), a(3)}}},
448452
{"sat", {"sat", {a(0)}}},
449453
{"fptoui_sat", {"fptoui.sat", {a(0)}}},
450454
{"fptosi_sat", {"fptosi.sat", {a(0)}}},
@@ -885,6 +889,34 @@ static void translateUnPackMask(CallInst &CI) {
885889
CI.replaceAllUsesWith(TransCI);
886890
}
887891

892+
// This function sets VCNamedBarrierCount attribute to set
893+
// the number of named barriers required by a kernel
894+
static void translateNbarrierInit(CallInst &CI) {
895+
auto *F = CI.getFunction();
896+
897+
auto *ArgV = CI.getArgOperand(0);
898+
assert(isa<ConstantInt>(ArgV) &&
899+
"integral constant expected for nbarrier count");
900+
901+
auto NewVal = cast<llvm::ConstantInt>(ArgV)->getZExtValue();
902+
assert(NewVal != 0 && "zero nbarrier count being requested");
903+
904+
if (llvm::MDNode *Node = getSLMSizeMDNode(F)) {
905+
if (llvm::Value *OldCount =
906+
getVal(Node->getOperand(genx::KernelMDOp::NBarrierCnt))) {
907+
assert(isa<llvm::ConstantInt>(OldCount) && "integer constant expected");
908+
llvm::Value *NewCount =
909+
llvm::ConstantInt::get(OldCount->getType(), NewVal);
910+
uint64_t OldVal = cast<llvm::ConstantInt>(OldCount)->getZExtValue();
911+
if (OldVal < NewVal)
912+
Node->replaceOperandWith(genx::KernelMDOp::NBarrierCnt,
913+
getMD(NewCount));
914+
}
915+
} else {
916+
llvm_unreachable("esimd_nbarrier_init can only be called by a kernel");
917+
}
918+
}
919+
888920
static bool translateVLoad(CallInst &CI, SmallPtrSet<Type *, 4> &GVTS) {
889921
if (GVTS.find(CI.getType()) != GVTS.end())
890922
return false;
@@ -1406,7 +1438,10 @@ void generateKernelMetadata(Module &M) {
14061438
getMD(llvm::ConstantInt::getNullValue(I32Ty)), // SLM size in bytes
14071439
getMD(llvm::ConstantInt::getNullValue(I32Ty)), // arg offsets
14081440
IOKinds,
1409-
ArgDescs};
1441+
ArgDescs,
1442+
getMD(llvm::ConstantInt::getNullValue(I32Ty)), // named barrier count
1443+
getMD(llvm::ConstantInt::getNullValue(I32Ty)) // regular barrier count
1444+
};
14101445

14111446
// Add this kernel to the root.
14121447
Kernels->addOperand(MDNode::get(Ctx, MDArgs));
@@ -1521,12 +1556,17 @@ size_t SYCLLowerESIMDPass::runOnFunction(Function &F,
15211556
// process ESIMD builtins that go through special handling instead of
15221557
// the translation procedure
15231558
// TODO FIXME slm_init should be made top-level __esimd_slm_init
1524-
if (Name.startswith("N2cl4sycl3ext5intel12experimental5esimd8slm_init")) {
1559+
if (Name.startswith("__esimd_slm_init")) {
15251560
// tag the kernel with meta-data SLMSize, and remove this builtin
15261561
translateSLMInit(*CI);
15271562
ToErase.push_back(CI);
15281563
continue;
15291564
}
1565+
if (Name.startswith("__esimd_nbarrier_init")) {
1566+
translateNbarrierInit(*CI);
1567+
ToErase.push_back(CI);
1568+
continue;
1569+
}
15301570
if (Name.startswith("__esimd_pack_mask")) {
15311571
translatePackMask(*CI);
15321572
ToErase.push_back(CI);

‎llvm/test/SYCLLowerIR/ESIMD/acc_ptr.ll

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -30,7 +30,7 @@ define weak_odr dso_local spir_kernel void @ESIMDKernel(i32 %_arg_, float addrsp
3030

3131
; CHECK: attributes #[[GENX_MAIN]] = { "CMGenxMain" "oclrt"="1" }
3232
; CHECK: !genx.kernels = !{![[GENX_KERNELS:[0-9]+]]}
33-
; CHECK: ![[GENX_KERNELS]] = !{void (i32, float addrspace(1)*, float addrspace(1)*, i32, float addrspace(1)*)* @ESIMDKernel, !"ESIMDKernel", ![[ARG_KINDS:[0-9]+]], i32 0, i32 0, ![[ARG_IO_KINDS:[0-9]+]], ![[ARG_DESCS:[0-9]+]]}
33+
; CHECK: ![[GENX_KERNELS]] = !{void (i32, float addrspace(1)*, float addrspace(1)*, i32, float addrspace(1)*)* @ESIMDKernel, !"ESIMDKernel", ![[ARG_KINDS:[0-9]+]], i32 0, i32 0, ![[ARG_IO_KINDS:[0-9]+]], ![[ARG_DESCS:[0-9]+]], i32 0, i32 0}
3434
; CHECK: ![[ARG_KINDS]] = !{i32 0, i32 2, i32 2, i32 0, i32 0}
3535
; CHECK: ![[ARG_IO_KINDS]] = !{i32 0, i32 0, i32 0, i32 0, i32 0}
3636
; CHECK: ![[ARG_DESCS]] = !{!"", !"buffer_t", !"buffer_t", !"", !"svmptr_t"}

‎llvm/test/SYCLLowerIR/ESIMD/lower_intrins.ll

Lines changed: 0 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -81,7 +81,6 @@ define dso_local spir_kernel void @FUNC_30() {
8181
; CHECK: define dso_local spir_kernel void @FUNC_30()
8282
call spir_func void @_ZN2cl4sycl3ext5intel12experimental5esimd8slm_initEj(i32 1023)
8383
ret void
84-
; CHECK-NEXT: ret void
8584
}
8685

8786
define dso_local spir_func <16 x i32> @FUNC_32() {
@@ -327,6 +326,5 @@ attributes #0 = { "genx_byte_offset"="192" "genx_volatile" }
327326
!genx.kernels = !{!0}
328327

329328
!0 = !{void ()* @"FUNC_30", !"FUNC_30", !1, i32 0, i32 0, !1, !2, i32 0, i32 0}
330-
; CHECK: !0 = !{void ()* @FUNC_30, !"FUNC_30", !1, i32 1023, i32 0, !1, !2, i32 0, i32 0}
331329
!1 = !{i32 0, i32 0}
332330
!2 = !{}

‎llvm/test/SYCLLowerIR/esimd_lower_nbarriers.ll

Lines changed: 21 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,21 @@
1+
; RUN: opt < %s -LowerESIMD -S | FileCheck %s
2+
3+
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"
4+
target triple = "spir64-unknown-unknown-sycldevice"
5+
6+
; Function Attrs: convergent norecurse mustprogress
7+
define dso_local spir_kernel void @_ZTSZ6calleriE12kernel_esimd() !sycl_explicit_simd !3 {
8+
entry:
9+
; CHECK: call void @llvm.genx.nbarrier(i8 0, i8 2, i8 0)
10+
call spir_func void @_Z16__esimd_nbarrierhhh(i8 zeroext 0, i8 zeroext 2, i8 zeroext 0)
11+
12+
; CHECK: call void @llvm.genx.raw.send.noresult.i1.v8i32(i32 0, i1 true, i32 3, i32 33554436, <8 x i32> <i32 0, i32 0, i32 67371008, i32 0, i32 0, i32 0, i32 0, i32 0>)
13+
call spir_func void @_Z32__esimd_raw_send_nbarrier_signalIjLi8EEvjjjN2cl4sycl5INTEL3gpu6detail11vector_typeIT_XT0_EE4typeEt(i32 0, i32 3, i32 33554436, <8 x i32> <i32 0, i32 0, i32 67371008, i32 0, i32 0, i32
14+
0, i32 0, i32 0>, i16 zeroext 1)
15+
16+
ret void
17+
}
18+
!3 = !{}
19+
20+
declare dso_local spir_func void @_Z16__esimd_nbarrierhhh(i8 zeroext, i8 zeroext, i8 zeroext) local_unnamed_addr #1
21+
declare dso_local spir_func void @_Z32__esimd_raw_send_nbarrier_signalIjLi8EEvjjjN2cl4sycl5INTEL3gpu6detail11vector_typeIT_XT0_EE4typeEt(i32, i32, i32, <8 x i32>, i16 zeroext)

‎llvm/test/tools/sycl-post-link/spec-constants/SYCL2020-struct-with-undef-padding.ll

Lines changed: 89 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,89 @@
1+
; RUN: sycl-post-link --spec-const=rt -S %s -o %t.files.table
2+
; RUN: FileCheck %s -input-file=%t.files_0.ll --check-prefix CHECK-IR
3+
; RUN: FileCheck %s -input-file=%t.files_0.prop --check-prefix CHECK-PROP
4+
;
5+
; This test is intended to check that SpecConstantsPass is able to handle the
6+
; situation where specialization constants with complex types such as structs
7+
; have an 'undef' value for padding in LLVM IR
8+
9+
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"
10+
target triple = "spir64-unknown-unknown"
11+
12+
%"class.cl::sycl::specialization_id" = type { %struct.coeff_str_aligned_t }
13+
%"class.cl::sycl::specialization_id.1" = type { %struct.coeff2_str_aligned_t }
14+
%struct.coeff_str_aligned_t = type { %"class.std::array", i64, [8 x i8] }
15+
%struct.coeff2_str_aligned_t = type { %"class.std::array", i64, [7 x i8], i8 }
16+
%"class.std::array" = type { [3 x float] }
17+
18+
$_ZTSZ4mainEUlN2cl4sycl14kernel_handlerEE_ = comdat any
19+
20+
@__usid_str = private unnamed_addr constant [32 x i8] c"ef880fa09cf7a9d7____ZL8coeff_id\00", align 1
21+
@_ZL8coeff_id = internal addrspace(1) constant %"class.cl::sycl::specialization_id" { %struct.coeff_str_aligned_t { %"class.std::array" zeroinitializer, i64 0, [8 x i8] undef } }, align 32
22+
@__usid_str.0 = private unnamed_addr constant [33 x i8] c"df991fa0adf9bad8____ZL8coeff_id2\00", align 1
23+
@_ZL8coeff_id2 = internal addrspace(1) constant %"class.cl::sycl::specialization_id.1" { %struct.coeff2_str_aligned_t { %"class.std::array" zeroinitializer, i64 0, [7 x i8] undef, i8 undef } }, align 32
24+
25+
; Function Attrs: convergent norecurse
26+
define weak_odr dso_local spir_kernel void @_ZTSZ4mainEUlN2cl4sycl14kernel_handlerEE_() local_unnamed_addr #0 comdat !kernel_arg_buffer_location !6 !sycl_kernel_omit_args !7 {
27+
%1 = alloca %struct.coeff_str_aligned_t, align 32
28+
%2 = addrspacecast %struct.coeff_str_aligned_t* %1 to %struct.coeff_str_aligned_t addrspace(4)*
29+
%3 = bitcast %struct.coeff_str_aligned_t* %1 to i8*
30+
call spir_func void @_Z40__sycl_getComposite2020SpecConstantValueI19coeff_str_aligned_tET_PKcPKvS5_(%struct.coeff_str_aligned_t addrspace(4)* sret(%struct.coeff_str_aligned_t) align 32 %2, i8 addrspace(4)* noundef addrspacecast (i8* getelementptr inbounds ([32 x i8], [32 x i8]* @__usid_str, i64 0, i64 0) to i8 addrspace(4)*), i8 addrspace(4)* noundef addrspacecast (i8 addrspace(1)* bitcast (%"class.cl::sycl::specialization_id" addrspace(1)* @_ZL8coeff_id to i8 addrspace(1)*) to i8 addrspace(4)*), i8 addrspace(4)* noundef null) #4
31+
; CHECK-IR: %[[#NS0:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID0:]], float 0.000000e+00)
32+
; CHECK-IR: %[[#NS1:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID1:]], float 0.000000e+00)
33+
; CHECK-IR: %[[#NS2:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID2:]], float 0.000000e+00)
34+
; CHECK-IR: %[[#NS3:]] = call [3 x float] @_Z29__spirv_SpecConstantCompositefff_RA3_f(float %[[#NS0]], float %[[#NS1]], float %[[#NS2]])
35+
; CHECK-IR: %[[#NS4:]] = call %"class.std::array" @"_Z29__spirv_SpecConstantCompositeA3_f_Rclass.std::array"([3 x float] %[[#NS3]])
36+
; CHECK-IR: %[[#NS5:]] = call i64 @_Z20__spirv_SpecConstantix(i32 [[#SCID3:]], i64 0)
37+
; CHECK-IR: %[[#NS6:]] = call %struct.coeff_str_aligned_t @"_Z29__spirv_SpecConstantCompositeclass.std::arrayxA8_a_Rstruct.coeff_str_aligned_t"(%"class.std::array" %[[#NS4]], i64 %[[#NS5]], [8 x i8] undef)
38+
39+
%4 = alloca %struct.coeff2_str_aligned_t, align 32
40+
%5 = addrspacecast %struct.coeff2_str_aligned_t* %4 to %struct.coeff2_str_aligned_t addrspace(4)*
41+
%6 = bitcast %struct.coeff2_str_aligned_t* %4 to i8*
42+
call spir_func void @_Z40__sycl_getComposite2020SpecConstantValueI19coeff2_str_aligned_tET_PKcPKvS5_(%struct.coeff2_str_aligned_t addrspace(4)* sret(%struct.coeff2_str_aligned_t) align 32 %5, i8 addrspace(4)* noundef addrspacecast (i8* getelementptr inbounds ([33 x i8], [33 x i8]* @__usid_str.0, i64 0, i64 0) to i8 addrspace(4)*), i8 addrspace(4)* noundef addrspacecast (i8 addrspace(1)* bitcast (%"class.cl::sycl::specialization_id.1" addrspace(1)* @_ZL8coeff_id2 to i8 addrspace(1)*) to i8 addrspace(4)*), i8 addrspace(4)* noundef null) #4
43+
; CHECK-IR: %[[#NS7:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID4:]], float 0.000000e+00)
44+
; CHECK-IR: %[[#NS8:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID5:]], float 0.000000e+00)
45+
; CHECK-IR: %[[#NS9:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID6:]], float 0.000000e+00)
46+
; CHECK-IR: %[[#NS10:]] = call [3 x float] @_Z29__spirv_SpecConstantCompositefff_RA3_f(float %[[#NS7]], float %[[#NS8]], float %[[#NS9]])
47+
; CHECK-IR: %[[#NS11:]] = call %"class.std::array" @"_Z29__spirv_SpecConstantCompositeA3_f_Rclass.std::array"([3 x float] %[[#NS10]])
48+
; CHECK-IR: %[[#NS12:]] = call i64 @_Z20__spirv_SpecConstantix(i32 [[#SCID7:]], i64 0)
49+
; CHECK-IR: %[[#NS13:]] = call %struct.coeff2_str_aligned_t @"_Z29__spirv_SpecConstantCompositeclass.std::arrayxA7_aa_Rstruct.coeff2_str_aligned_t"(%"class.std::array" %[[#NS11]], i64 %[[#NS12]], [7 x i8] undef, i8 undef)
50+
51+
ret void
52+
}
53+
; Function Attrs: convergent
54+
declare dso_local spir_func void @_Z40__sycl_getComposite2020SpecConstantValueI19coeff_str_aligned_tET_PKcPKvS5_(%struct.coeff_str_aligned_t addrspace(4)* sret(%struct.coeff_str_aligned_t) align 32, i8 addrspace(4)* noundef, i8 addrspace(4)* noundef, i8 addrspace(4)* noundef) local_unnamed_addr #2
55+
56+
declare dso_local spir_func void @_Z40__sycl_getComposite2020SpecConstantValueI19coeff2_str_aligned_tET_PKcPKvS5_(%struct.coeff2_str_aligned_t addrspace(4)* sret(%struct.coeff2_str_aligned_t) align 32, i8 addrspace(4)* noundef, i8 addrspace(4)* noundef, i8 addrspace(4)* noundef) local_unnamed_addr #2
57+
58+
attributes #0 = { convergent norecurse "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="spec-constant-test.cpp" "uniform-work-group-size"="true" }
59+
attributes #1 = { argmemonly mustprogress nofree nosync nounwind willreturn }
60+
attributes #2 = { convergent "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
61+
attributes #3 = { nounwind }
62+
attributes #4 = { convergent }
63+
64+
!llvm.dependent-libraries = !{!0}
65+
!llvm.module.flags = !{!1, !2}
66+
!opencl.spir.version = !{!3}
67+
!spirv.Source = !{!4}
68+
!llvm.ident = !{!5}
69+
; CHECK-IR: !sycl.specialization-constants = !{![[#MN0:]], ![[#MN1:]]}
70+
; CHECK-IR: !sycl.specialization-constants-default-values = !{![[#MN2:]], ![[#MN3:]]}
71+
72+
!0 = !{!"libcpmt"}
73+
!1 = !{i32 1, !"wchar_size", i32 2}
74+
!2 = !{i32 7, !"frame-pointer", i32 2}
75+
!3 = !{i32 1, i32 2}
76+
!4 = !{i32 4, i32 100000}
77+
!5 = !{!"clang version 14.0.0"}
78+
!6 = !{i32 -1}
79+
!7 = !{i1 true}
80+
; CHECK-IR: ![[#MN0]] = !{!"ef880fa09cf7a9d7____ZL8coeff_id", i32 0, i32 0, i32 4, i32 1, i32 4, i32 4, i32 2, i32 8, i32 4, i32 3, i32 16, i32 8, i32 -1, i32 24, i32 8}
81+
; CHECK-IR: ![[#MN1]] = !{!"df991fa0adf9bad8____ZL8coeff_id2", i32 5, i32 0, i32 4, i32 6, i32 4, i32 4, i32 7, i32 8, i32 4, i32 8, i32 16, i32 8, i32 -1, i32 31, i32 1}
82+
; CHECK-IR: ![[#MN2]] = !{%struct.coeff_str_aligned_t { %"class.std::array" zeroinitializer, i64 0, [8 x i8] undef }}
83+
; CHECK-IR: ![[#MN3]] = !{%struct.coeff2_str_aligned_t { %"class.std::array" zeroinitializer, i64 0, [7 x i8] undef, i8 undef }}
84+
85+
; CHECK-PROP: [SYCL/specialization constants]
86+
; CHECK-PROP-NEXT: ef880fa09cf7a9d7____ZL8coeff_id=2|
87+
88+
; CHECK-PROP: [SYCL/specialization constants default values]
89+
; CHECK-PROP-NEXT: all=2|

‎llvm/test/tools/sycl-post-link/sycl-esimd/nbarriers-metadata.ll

Lines changed: 19 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,19 @@
1+
; RUN: sycl-post-link -split-esimd -lower-esimd -S %s -o %t.table
2+
; RUN: FileCheck %s -input-file=%t_esimd_0.ll
3+
4+
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"
5+
target triple = "spir64-unknown-unknown-sycldevice"
6+
7+
; Function Attrs: convergent norecurse mustprogress
8+
define dso_local spir_kernel void @_ZTSZ6calleriE12kernel_esimd() #0 !sycl_explicit_simd !3 {
9+
entry:
10+
tail call spir_func void @_Z21__esimd_nbarrier_inith(i8 zeroext 7)
11+
ret void
12+
}
13+
14+
!3 = !{}
15+
16+
declare dso_local spir_func void @_Z21__esimd_nbarrier_inith(i8 zeroext)
17+
; CHECK: attributes #0 = { {{.*}}"VCNamedBarrierCount"="7"{{.*}} }
18+
19+
attributes #0 = { "sycl-module-id"="a.cpp" }

‎llvm/tools/sycl-post-link/CompileTimePropertiesPass.h

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -63,7 +63,7 @@ inline bool hasProperty(const Attribute &Attr) {
6363
template <typename Int> Int getAttributeAsInteger(const Attribute &Attr) {
6464
assert(Attr.isStringAttribute() &&
6565
"The attribute Attr must be a string attribute");
66-
Int Value;
66+
Int Value = 0;
6767
bool Error = Attr.getValueAsString().getAsInteger(10, Value);
6868
assert(!Error && "The attribute's value is not a number");
6969
(void)Error;

0 commit comments

Comments
 (0)