Skip to content

[SYCL] Add support to propagate compile flags to device backend compiler #8763

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
47 commits
Select commit Hold shift + click to select a range
fb3bebd
Add frontend changes
asudarsa Mar 20, 2023
d3a9e24
Adding support to propagate frontend compile time options to backend
asudarsa Mar 23, 2023
395d550
Minor changes
asudarsa Mar 24, 2023
0467da0
Modify order of split modules in the sycl-post-link table
asudarsa Mar 24, 2023
1cc4d14
Add documentation
asudarsa Mar 24, 2023
5c529d2
Merge remote-tracking branch 'intel_llvm/sycl' into asudarsa/add_supp…
asudarsa Mar 24, 2023
570d720
Fix formatting issues
asudarsa Mar 24, 2023
992614e
Add plugin API to get backend option
asudarsa Mar 27, 2023
e2e678a
Address review comments on tests
asudarsa Mar 27, 2023
d89738e
Address review comments and other changes to documentation
asudarsa Mar 27, 2023
67ac752
Minor correction
asudarsa Mar 27, 2023
87a6f14
Move attribute addition to a separate pass called in backend utils pa…
asudarsa Mar 29, 2023
97fdae7
Address review comments on documentation
asudarsa Mar 29, 2023
1f2015e
minor typo
asudarsa Mar 29, 2023
2cc0bca
Removed check for a specific attribute id
asudarsa Mar 29, 2023
b77fc0f
Move test from llvm-test-suite
asudarsa Mar 29, 2023
aa878cc
Address review comments
asudarsa Mar 29, 2023
6e52273
Remove warning message
asudarsa Mar 29, 2023
d427727
Merge remote-tracking branch 'intel_llvm/sycl' into asudarsa/add_supp…
asudarsa Apr 1, 2023
41984be
Update design document and plugin implementation
asudarsa Apr 4, 2023
a34e009
document correction
asudarsa Apr 4, 2023
e14db04
document correction
asudarsa Apr 4, 2023
6359987
document correction
asudarsa Apr 4, 2023
b021163
Addressing code review changes
asudarsa Apr 5, 2023
f2d98a2
Address more review comments
asudarsa Apr 5, 2023
22f034b
Change char strs in plugin to static const char strs
asudarsa Apr 5, 2023
ce810ee
Address SYCL RT review comments
asudarsa Apr 6, 2023
a9c01b7
remove unneeded checks in tests
asudarsa Apr 6, 2023
15f4abb
Merge issues
asudarsa Apr 6, 2023
14e882a
format issue fixed
asudarsa Apr 6, 2023
476a2f0
Modify plugin changes based on user reviews
asudarsa Apr 7, 2023
1cfa88b
Merge remote-tracking branch 'intel_llvm/sycl' into asudarsa/add_supp…
asudarsa Apr 7, 2023
ad50e96
fixing unrelated format change to make tests proceed
asudarsa Apr 7, 2023
36dda04
Fix unused variable warning
asudarsa Apr 7, 2023
7f91227
Fix test fails
asudarsa Apr 7, 2023
2f50ba0
Turning off option passing for ESIMD images
asudarsa Apr 7, 2023
e25105c
More improvements based on review comments
asudarsa Apr 7, 2023
8ffca4b
Fix format issues
asudarsa Apr 7, 2023
f1fadc4
More review comments addressed
asudarsa Apr 7, 2023
ef009ff
Minor changes in plugin code
asudarsa Apr 7, 2023
b5b1b75
Removed unused variable and extra parentheses
asudarsa Apr 7, 2023
3937a2a
Minor format issue
asudarsa Apr 7, 2023
d038d4e
Minor typos fixed
asudarsa Apr 7, 2023
581d22a
Merge remote-tracking branch 'intel_llvm/sycl' into asudarsa/add_supp…
asudarsa Apr 7, 2023
821a549
Merge remote-tracking branch 'intel_llvm/sycl' into asudarsa/add_supp…
asudarsa Apr 10, 2023
c26c2fd
Merge remote-tracking branch 'intel_llvm/sycl' into asudarsa/add_supp…
asudarsa Apr 12, 2023
a1ce60f
Minor doc changes
asudarsa Apr 13, 2023
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
6 changes: 5 additions & 1 deletion clang/lib/CodeGen/BackendUtil.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -48,6 +48,7 @@
#include "llvm/SYCLLowerIR/ESIMD/ESIMDVerifier.h"
#include "llvm/SYCLLowerIR/LowerWGLocalMemory.h"
#include "llvm/SYCLLowerIR/MutatePrintfAddrspace.h"
#include "llvm/SYCLLowerIR/SYCLAddOptLevelAttribute.h"
#include "llvm/SYCLLowerIR/SYCLPropagateAspectsUsage.h"
#include "llvm/Support/BuryPointer.h"
#include "llvm/Support/CommandLine.h"
Expand All @@ -60,8 +61,8 @@
#include "llvm/Support/raw_ostream.h"
#include "llvm/Target/TargetMachine.h"
#include "llvm/Target/TargetOptions.h"
#include "llvm/Transforms/IPO/DeadArgumentElimination.h"
#include "llvm/TargetParser/Triple.h"
#include "llvm/Transforms/IPO/DeadArgumentElimination.h"
#include "llvm/Transforms/IPO/LowerTypeTests.h"
#include "llvm/Transforms/IPO/ThinLTOBitcodeWriter.h"
#include "llvm/Transforms/InstCombine/InstCombine.h"
Expand Down Expand Up @@ -1045,6 +1046,9 @@ void EmitAssemblyHelper::RunOptimizationPipeline(
MPM.addPass(SYCLPropagateAspectsUsagePass(/*ExcludeAspects=*/{},
/*ValidateAspects=*/false));

// Add attribute corresponding to optimization level.
MPM.addPass(SYCLAddOptLevelAttributePass(CodeGenOpts.OptimizationLevel));

// Add SPIRITTAnnotations pass to the pass manager if
// -fsycl-instrument-device-code option was passed. This option can be
// used only with spir triple.
Expand Down
18 changes: 18 additions & 0 deletions clang/test/CodeGenSYCL/sycl-add-opt-level-attrib.cpp
Original file line number Diff line number Diff line change
@@ -0,0 +1,18 @@
// RUN: %clangxx %s -O0 -S -o %t.ll -fsycl-device-only
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Do we need an additional test with explicit __attribute__((optnone)) on the kernel/device function?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

My changes do not interact with this attribute. The backend should be responsible to decide whether the attribute or the option takes precedence.

Thanks

// RUN: FileCheck %s --input-file %t.ll -check-prefixes=CHECK-IR
// CHECK-IR: define {{.*}} spir_kernel void @{{.*}}main{{.*}}sycl{{.*}}handler{{.*}}() #[[ATTR:[0-9]+]]
// CHECK-IR: attributes #[[ATTR]] = { {{.*}} "sycl-optlevel"="0" {{.*}}}

// This test checks adding of the attribute 'sycl-optlevel'
// by the clang front-end

#include <sycl/sycl.hpp>

int main() {
sycl::queue q;
q.submit([&](sycl::handler &h) {
h.single_task([=]() {});
});
return 0;
}

33 changes: 33 additions & 0 deletions llvm/include/llvm/SYCLLowerIR/SYCLAddOptLevelAttribute.h
Original file line number Diff line number Diff line change
@@ -0,0 +1,33 @@
//===----- SYCLAddOptLevelAttribute.h - SYCLAddOptLevelAttribute Pass -----===//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
//
//===----------------------------------------------------------------------===//
//
// Pass adds 'sycl-optlevel' function attribute based on optimization level
// passed in.
//
//===----------------------------------------------------------------------===//
//
#ifndef LLVM_SYCL_ADD_OPT_LEVEL_ATTRIBUTE_H
#define LLVM_SYCL_ADD_OPT_LEVEL_ATTRIBUTE_H

#include "llvm/IR/PassManager.h"

namespace llvm {

class SYCLAddOptLevelAttributePass
: public PassInfoMixin<SYCLAddOptLevelAttributePass> {
public:
SYCLAddOptLevelAttributePass(int OptLevel = -1) : OptLevel{OptLevel} {};
PreservedAnalyses run(Module &M, ModuleAnalysisManager &);

private:
int OptLevel;
};

} // namespace llvm

#endif // LLVM_SYCL_ADD_OPT_LEVEL_ATTRIBUTE_H
1 change: 1 addition & 0 deletions llvm/lib/Passes/PassBuilder.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -90,6 +90,7 @@
#include "llvm/SYCLLowerIR/LowerWGLocalMemory.h"
#include "llvm/SYCLLowerIR/LowerWGScope.h"
#include "llvm/SYCLLowerIR/MutatePrintfAddrspace.h"
#include "llvm/SYCLLowerIR/SYCLAddOptLevelAttribute.h"
#include "llvm/SYCLLowerIR/SYCLPropagateAspectsUsage.h"
#include "llvm/Support/CommandLine.h"
#include "llvm/Support/Debug.h"
Expand Down
1 change: 1 addition & 0 deletions llvm/lib/Passes/PassRegistry.def
Original file line number Diff line number Diff line change
Expand Up @@ -140,6 +140,7 @@ MODULE_PASS("deadargelim-sycl", DeadArgumentEliminationSYCLPass())
MODULE_PASS("sycllowerwglocalmemory", SYCLLowerWGLocalMemoryPass())
MODULE_PASS("lower-esimd-kernel-attrs", SYCLFixupESIMDKernelWrapperMDPass())
MODULE_PASS("sycl-propagate-aspects-usage", SYCLPropagateAspectsUsagePass())
MODULE_PASS("sycl-add-opt-level-attribute", SYCLAddOptLevelAttributePass())
MODULE_PASS("compile-time-properties", CompileTimePropertiesPass())
#undef MODULE_PASS

Expand Down
1 change: 1 addition & 0 deletions llvm/lib/SYCLLowerIR/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -63,6 +63,7 @@ add_llvm_component_library(LLVMSYCLLowerIR
LowerWGLocalMemory.cpp
LowerWGScope.cpp
MutatePrintfAddrspace.cpp
SYCLAddOptLevelAttribute.cpp
SYCLPropagateAspectsUsage.cpp
SYCLUtils.cpp

Expand Down
30 changes: 30 additions & 0 deletions llvm/lib/SYCLLowerIR/SYCLAddOptLevelAttribute.cpp
Original file line number Diff line number Diff line change
@@ -0,0 +1,30 @@
//===---- SYCLAddOptLevelAttribute.cpp - SYCLAddOptLevelAttribute Pass ---===//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
//
//===---------------------------------------------------------------------===//
//
// Pass adds 'sycl-optlevel' function attribute based on optimization level
// passed in.
//===---------------------------------------------------------------------===//

#include "llvm/SYCLLowerIR/SYCLAddOptLevelAttribute.h"

#include "llvm/IR/Module.h"

using namespace llvm;

PreservedAnalyses
SYCLAddOptLevelAttributePass::run(Module &M, ModuleAnalysisManager &MAM) {
// Here, we add a function attribute 'sycl-optlevel' to store the
// optimization level.
assert(OptLevel >= 0 && "Invalid optimization level!");
for (Function &F : M.functions()) {
if (F.isDeclaration())
continue;
F.addFnAttr("sycl-optlevel", std::to_string(OptLevel));
}
return PreservedAnalyses::all();
}
Original file line number Diff line number Diff line change
Expand Up @@ -10,29 +10,29 @@
; RUN: sycl-post-link -split=auto -symbols -S < %s -o %t.table
; RUN: FileCheck %s -input-file=%t_0.ll --check-prefixes CHECK-M0-IR \
; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1
; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-M1-IR \
; RUN: FileCheck %s -input-file=%t_2.ll --check-prefixes CHECK-M1-IR \
Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

These test updates were triggered by the fact that the order of the entries in the .table file changed due to the addition of the new entry in the list optional kernel features. A better way to NOT depend on the order of the entries is required in the long run.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

FYI: that unstable ordering is going to be fixed in #8833. That PR will also allow to simplify your changes to device code split down to a single line of code

; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1
; RUN: FileCheck %s -input-file=%t_2.ll --check-prefixes CHECK-M2-IR \
; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-M2-IR \
; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1
; RUN: FileCheck %s -input-file=%t_0.sym --check-prefixes CHECK-M0-SYMS \
; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1
; RUN: FileCheck %s -input-file=%t_1.sym --check-prefixes CHECK-M1-SYMS \
; RUN: FileCheck %s -input-file=%t_2.sym --check-prefixes CHECK-M1-SYMS \
; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1
; RUN: FileCheck %s -input-file=%t_2.sym --check-prefixes CHECK-M2-SYMS \
; RUN: FileCheck %s -input-file=%t_1.sym --check-prefixes CHECK-M2-SYMS \
; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1

; RUN: sycl-post-link -split=source -symbols -S < %s -o %t.table
; RUN: FileCheck %s -input-file=%t_0.ll --check-prefixes CHECK-M0-IR \
; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1
; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-M1-IR \
; RUN: FileCheck %s -input-file=%t_2.ll --check-prefixes CHECK-M1-IR \
; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1
; RUN: FileCheck %s -input-file=%t_2.ll --check-prefixes CHECK-M2-IR \
; RUN: FileCheck %s -input-file=%t_1.ll --check-prefixes CHECK-M2-IR \
; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1
; RUN: FileCheck %s -input-file=%t_0.sym --check-prefixes CHECK-M0-SYMS \
; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1
; RUN: FileCheck %s -input-file=%t_1.sym --check-prefixes CHECK-M1-SYMS \
; RUN: FileCheck %s -input-file=%t_2.sym --check-prefixes CHECK-M1-SYMS \
; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1
; RUN: FileCheck %s -input-file=%t_2.sym --check-prefixes CHECK-M2-SYMS \
; RUN: FileCheck %s -input-file=%t_1.sym --check-prefixes CHECK-M2-SYMS \
; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1

; RUN: sycl-post-link -split=kernel -symbols -S < %s -o %t.table
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -4,14 +4,14 @@
; RUN: sycl-post-link -split=auto -symbols -S < %s -o %t.table
; RUN: FileCheck %s -input-file=%t.table --check-prefix CHECK-TABLE
;
; RUN: FileCheck %s -input-file=%t_0.sym --check-prefix CHECK-M0-SYMS \
; RUN: FileCheck %s -input-file=%t_1.sym --check-prefix CHECK-M0-SYMS \
; RUN: --implicit-check-not kernel0 --implicit-check-not kernel1 \
; RUN: --implicit-check-not kernel2
;
; RUN: FileCheck %s -input-file=%t_1.sym --check-prefix CHECK-M2-SYMS \
; RUN: FileCheck %s -input-file=%t_2.sym --check-prefix CHECK-M2-SYMS \
; RUN: --implicit-check-not kernel0 --implicit-check-not kernel3
;
; RUN: FileCheck %s -input-file=%t_2.sym --check-prefix CHECK-M1-SYMS \
; RUN: FileCheck %s -input-file=%t_0.sym --check-prefix CHECK-M1-SYMS \
; RUN: --implicit-check-not kernel1 --implicit-check-not kernel2 \
; RUN: --implicit-check-not kernel3

Expand Down
16 changes: 8 additions & 8 deletions llvm/test/tools/sycl-post-link/sycl-esimd-large-grf.ll
Original file line number Diff line number Diff line change
Expand Up @@ -9,16 +9,16 @@

; RUN: sycl-post-link -split=source -symbols -split-esimd -lower-esimd -S < %s -o %t.table
; RUN: FileCheck %s -input-file=%t.table
; RUN: FileCheck %s -input-file=%t_esimd_large_grf_0.ll --check-prefixes CHECK-ESIMD-LargeGRF-IR --implicit-check-not='RegisterAllocMode'
; RUN: FileCheck %s -input-file=%t_esimd_large_grf_0.prop --check-prefixes CHECK-ESIMD-LargeGRF-PROP
; RUN: FileCheck %s -input-file=%t_1.sym --check-prefixes CHECK-SYCL-SYM
; RUN: FileCheck %s -input-file=%t_esimd_1.sym --check-prefixes CHECK-ESIMD-SYM
; RUN: FileCheck %s -input-file=%t_esimd_large_grf_0.sym --check-prefixes CHECK-ESIMD-LargeGRF-SYM
; RUN: FileCheck %s -input-file=%t_esimd_large_grf_1.ll --check-prefixes CHECK-ESIMD-LargeGRF-IR --implicit-check-not='RegisterAllocMode'
; RUN: FileCheck %s -input-file=%t_esimd_large_grf_1.prop --check-prefixes CHECK-ESIMD-LargeGRF-PROP
; RUN: FileCheck %s -input-file=%t_0.sym --check-prefixes CHECK-SYCL-SYM
; RUN: FileCheck %s -input-file=%t_esimd_0.sym --check-prefixes CHECK-ESIMD-SYM
; RUN: FileCheck %s -input-file=%t_esimd_large_grf_1.sym --check-prefixes CHECK-ESIMD-LargeGRF-SYM

; CHECK: [Code|Properties|Symbols]
; CHECK: {{.*}}_0.ll|{{.*}}_0.prop|{{.*}}_0.sym
; CHECK: {{.*}}esimd-large-grf.ll.tmp_1.ll|{{.*}}esimd-large-grf.ll.tmp_1.prop|{{.*}}esimd-large-grf.ll.tmp_1.sym
; CHECK: {{.*}}esimd-large-grf.ll.tmp_esimd_1.ll|{{.*}}esimd-large-grf.ll.tmp_esimd_1.prop|{{.*}}esimd-large-grf.ll.tmp_esimd_1.sym
; CHECK: {{.*}}esimd-large-grf.ll.tmp_0.ll|{{.*}}esimd-large-grf.ll.tmp_0.prop|{{.*}}esimd-large-grf.ll.tmp_0.sym
; CHECK: {{.*}}esimd-large-grf.ll.tmp_esimd_0.ll|{{.*}}esimd-large-grf.ll.tmp_esimd_0.prop|{{.*}}esimd-large-grf.ll.tmp_esimd_0.sym
; CHECK: {{.*}}_1.ll|{{.*}}_1.prop|{{.*}}_1.sym

; CHECK-ESIMD-LargeGRF-PROP: isEsimdImage=1|1
; CHECK-ESIMD-LargeGRF-PROP: isLargeGRF=1|1
Expand Down
12 changes: 6 additions & 6 deletions llvm/test/tools/sycl-post-link/sycl-large-grf.ll
Original file line number Diff line number Diff line change
Expand Up @@ -9,14 +9,14 @@

; RUN: sycl-post-link -split=source -symbols -split-esimd -lower-esimd -S < %s -o %t.table
; RUN: FileCheck %s -input-file=%t.table
; RUN: FileCheck %s -input-file=%t_large_grf_0.ll --check-prefixes CHECK-LARGE-GRF-IR
; RUN: FileCheck %s -input-file=%t_large_grf_0.prop --check-prefixes CHECK-LARGE-GRF-PROP
; RUN: FileCheck %s -input-file=%t_1.sym --check-prefixes CHECK-SYCL-SYM
; RUN: FileCheck %s -input-file=%t_large_grf_0.sym --check-prefixes CHECK-LARGE-GRF-SYM
; RUN: FileCheck %s -input-file=%t_large_grf_1.ll --check-prefixes CHECK-LARGE-GRF-IR
; RUN: FileCheck %s -input-file=%t_large_grf_1.prop --check-prefixes CHECK-LARGE-GRF-PROP
; RUN: FileCheck %s -input-file=%t_0.sym --check-prefixes CHECK-SYCL-SYM
; RUN: FileCheck %s -input-file=%t_large_grf_1.sym --check-prefixes CHECK-LARGE-GRF-SYM

; CHECK: [Code|Properties|Symbols]
; CHECK: {{.*}}_0.ll|{{.*}}_0.prop|{{.*}}_0.sym
; CHECK: {{.*}}-large-grf.ll.tmp_1.ll|{{.*}}-large-grf.ll.tmp_1.prop|{{.*}}-large-grf.ll.tmp_1.sym
; CHECK: {{.*}}-large-grf.ll.tmp_0.ll|{{.*}}-large-grf.ll.tmp_0.prop|{{.*}}-large-grf.ll.tmp_0.sym
; CHECK: {{.*}}_1.ll|{{.*}}_1.prop|{{.*}}_1.sym

; CHECK-LARGE-GRF-PROP: isLargeGRF=1|1

Expand Down
46 changes: 46 additions & 0 deletions llvm/test/tools/sycl-post-link/sycl-opt-level.ll
Original file line number Diff line number Diff line change
@@ -0,0 +1,46 @@
; This test checks parsing of the attribute 'sycl-optlevel'
; by the sycl-post-link-tool:
; In addition to splitting requested by user, the kernels are also split based
; on their optimization levels.
; sycl-post-link adds 'optLevel' property to the device binary

; RUN: sycl-post-link -split=source -symbols -S < %s -o %t.table
; RUN: FileCheck %s -input-file=%t.table
; RUN: FileCheck %s -input-file=%t_0.prop --check-prefixes CHECK-OPT-LEVEL-PROP-0
; RUN: FileCheck %s -input-file=%t_1.prop --check-prefixes CHECK-OPT-LEVEL-PROP-1

; CHECK: [Code|Properties|Symbols]
; CHECK: {{.*}}_0.ll|{{.*}}_0.prop|{{.*}}_0.sym
; CHECK: {{.*}}_1.ll|{{.*}}_1.prop|{{.*}}_1.sym

; CHECK-OPT-LEVEL-PROP-0: optLevel=1|0
; CHECK-OPT-LEVEL-PROP-1: optLevel=1|2

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"
target triple = "spir64-unknown-unknown"

define dso_local spir_func noundef i32 @_Z3fooii(i32 noundef %a, i32 noundef %b) local_unnamed_addr #0 {
entry:
%sub = sub nsw i32 %a, %b
ret i32 %sub
}

define dso_local spir_func noundef i32 @_Z3booii(i32 noundef %a, i32 noundef %b) #1 {
entry:
%retval = alloca i32, align 4
%a.addr = alloca i32, align 4
%b.addr = alloca i32, align 4
%retval.ascast = addrspacecast i32* %retval to i32 addrspace(4)*
%a.addr.ascast = addrspacecast i32* %a.addr to i32 addrspace(4)*
%b.addr.ascast = addrspacecast i32* %b.addr to i32 addrspace(4)*
store i32 %a, i32 addrspace(4)* %a.addr.ascast, align 4
store i32 %b, i32 addrspace(4)* %b.addr.ascast, align 4
%0 = load i32, i32 addrspace(4)* %a.addr.ascast, align 4
%1 = load i32, i32 addrspace(4)* %b.addr.ascast, align 4
%add = add nsw i32 %0, %1
ret i32 %add
}

attributes #0 = { mustprogress nofree norecurse nosync nounwind willreturn memory(none) "sycl-module-id"="test3.cpp" "sycl-optlevel"="2" }
attributes #1 = { convergent mustprogress noinline norecurse nounwind optnone "sycl-module-id"="test2.cpp" "sycl-optlevel"="0" }

24 changes: 20 additions & 4 deletions llvm/tools/sycl-post-link/ModuleSplitter.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -42,6 +42,7 @@ constexpr char GLOBAL_SCOPE_NAME[] = "<GLOBAL>";
constexpr char SYCL_SCOPE_NAME[] = "<SYCL>";
constexpr char ESIMD_SCOPE_NAME[] = "<ESIMD>";
constexpr char ESIMD_MARKER_MD[] = "sycl_explicit_simd";
constexpr char ATTR_OPT_LEVEL[] = "sycl-optlevel";

bool hasIndirectFunctionsOrCalls(const Module &M) {
for (const auto &F : M.functions()) {
Expand Down Expand Up @@ -674,7 +675,8 @@ void ModuleDesc::dump() const {
llvm::errs() << " ESIMD:" << toString(EntryPoints.Props.HasESIMD)
<< ", SpecConstMet:" << (Props.SpecConstsMet ? "YES" : "NO")
<< ", LargeGRF:"
<< (EntryPoints.Props.UsesLargeGRF ? "YES" : "NO") << "\n";
<< (EntryPoints.Props.UsesLargeGRF ? "YES" : "NO")
<< ", OptLevel:" << EntryPoints.getOptLevel() << "\n";
dumpEntryPoints(entries(), EntryPoints.GroupId.c_str(), 1);
llvm::errs() << "}\n";
}
Expand Down Expand Up @@ -713,6 +715,7 @@ namespace {
struct UsedOptionalFeatures {
SmallVector<int, 4> Aspects;
bool UsesLargeGRF = false;
int OptLevel = -1;
SmallVector<int, 3> ReqdWorkGroupSize;
// TODO: extend this further with reqd-sub-group-size and other properties

Expand All @@ -735,6 +738,12 @@ struct UsedOptionalFeatures {
if (F->hasFnAttribute(::sycl::kernel_props::ATTR_LARGE_GRF))
UsesLargeGRF = true;

if (F->hasFnAttribute(ATTR_OPT_LEVEL))
if (F->getFnAttribute(ATTR_OPT_LEVEL)
.getValueAsString()
.getAsInteger(10, OptLevel))
OptLevel = -1;

if (const MDNode *MDN = F->getMetadata("reqd_work_group_size")) {
size_t NumOperands = MDN->getNumOperands();
assert(NumOperands >= 1 && NumOperands <= 3 &&
Expand All @@ -750,8 +759,9 @@ struct UsedOptionalFeatures {
llvm::hash_code LargeGRFHash = llvm::hash_value(UsesLargeGRF);
llvm::hash_code ReqdWorkGroupSizeHash = llvm::hash_combine_range(
ReqdWorkGroupSize.begin(), ReqdWorkGroupSize.end());
Hash = static_cast<unsigned>(
llvm::hash_combine(AspectsHash, LargeGRFHash, ReqdWorkGroupSizeHash));
llvm::hash_code OptLevelHash = llvm::hash_value(OptLevel);
Hash = static_cast<unsigned>(llvm::hash_combine(
AspectsHash, LargeGRFHash, ReqdWorkGroupSizeHash, OptLevelHash));
}

std::string generateModuleName(StringRef BaseName) const {
Expand All @@ -773,6 +783,9 @@ struct UsedOptionalFeatures {
if (UsesLargeGRF)
Ret += "-large-grf";

if (OptLevel != -1)
Ret += "-O" + std::to_string(OptLevel);

return Ret;
}

Expand Down Expand Up @@ -808,7 +821,8 @@ struct UsedOptionalFeatures {
return false;
}

return IsEmpty == Other.IsEmpty && UsesLargeGRF == Other.UsesLargeGRF;
return IsEmpty == Other.IsEmpty && UsesLargeGRF == Other.UsesLargeGRF &&
OptLevel == Other.OptLevel;
}

unsigned hash() const { return static_cast<unsigned>(Hash); }
Expand Down Expand Up @@ -869,6 +883,8 @@ getSplitterByOptionalFeatures(ModuleDesc &&MD,
// Propagate LargeGRF flag to entry points group
if (Features.UsesLargeGRF)
MDProps.UsesLargeGRF = true;
if (Features.OptLevel != -1)
MDProps.OptLevel = Features.OptLevel;
Groups.emplace_back(
Features.generateModuleName(MD.getEntryPointGroup().GroupId),
std::move(EntryPoints), MDProps);
Expand Down
Loading