Skip to content

Commit cb3022f

Browse files
committed
Handle SPIR AOT
1 parent 02fff2b commit cb3022f

File tree

2 files changed

+31
-0
lines changed

2 files changed

+31
-0
lines changed
Lines changed: 24 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,24 @@
1+
; RUN: sycl-post-link -spec-const=emulation %s 2>&1 | FileCheck %s
2+
3+
; This test checks the `-spec-const` pass on SPIR-V targets and emulation mode,
4+
; i.e., on AOT SPIR-V targets. In this scenario, 'llvm.sycl.alloca' intrinsics
5+
; must be left unmodified.
6+
7+
; Note that coming from clang this case should never be reached.
8+
9+
; CHECK: sycl-post-link NOTE: no modifications to the input LLVM IR have been made
10+
11+
target triple = "spir64_x86_64"
12+
13+
%"class.sycl::_V1::specialization_id" = type { i64 }
14+
15+
@size_i64 = addrspace(1) constant %"class.sycl::_V1::specialization_id" { i64 10 }, align 8
16+
17+
@size_i64_stable_name = private unnamed_addr constant [36 x i8] c"_ZTS14name_generatorIL_Z8size_i64EE\00", align 1
18+
19+
define dso_local void @private_alloca() {
20+
call ptr @llvm.sycl.alloca.p0.p4.p4.p4.f64(ptr addrspace(4) addrspacecast (ptr @size_i64_stable_name to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr addrspace(1) @size_i64 to ptr addrspace(4)), ptr addrspace(4) null, double 0.000000e+00, i64 8)
21+
ret void
22+
}
23+
24+
declare ptr @llvm.sycl.alloca.p0.p4.p4.p4.f32(ptr addrspace(4), ptr addrspace(4), ptr addrspace(4), float, i64)

llvm/tools/sycl-post-link/SpecConstants.cpp

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -20,6 +20,7 @@
2020
#include "llvm/IR/Instructions.h"
2121
#include "llvm/IR/IntrinsicInst.h"
2222
#include "llvm/IR/Operator.h"
23+
#include "llvm/TargetParser/Triple.h"
2324

2425
#include <vector>
2526

@@ -815,12 +816,18 @@ PreservedAnalyses SpecConstantsPass::run(Module &M,
815816
// intrinsic to find its calls and lower them depending on the HandlingMode.
816817
bool IRModified = false;
817818
LLVMContext &Ctx = M.getContext();
819+
bool IsSPIREmulated =
820+
Triple(M.getTargetTriple()).isSPIR() && Mode == HandlingMode::emulation;
818821
for (Function &F : M) {
819822
if (!F.isDeclaration())
820823
continue;
821824

822825
const bool IsSYCLAlloca = F.getIntrinsicID() == Intrinsic::sycl_alloca;
823826

827+
// 'llvm.sycl.alloca' is not supported in emulation mode on SPIR-V targets.
828+
if (IsSPIREmulated && IsSYCLAlloca)
829+
continue;
830+
824831
if (!F.getName().starts_with(SYCL_GET_SCALAR_2020_SPEC_CONST_VAL) &&
825832
!F.getName().starts_with(SYCL_GET_COMPOSITE_2020_SPEC_CONST_VAL) &&
826833
!IsSYCLAlloca)

0 commit comments

Comments
 (0)