-
Notifications
You must be signed in to change notification settings - Fork 730
Commit
This commit does not belong to any branch on this repository, and may belong to a fork outside of the repository.
[SYCL] Reimplemented -f[no]sycl-early-optimizations flag (#7701)
Reimplemented the `-f[no]sycl-early-optimizations` compiler flag to separate it from the meaning of `-disable-llvm-passes` for more fidelity. This required a change to its definition, setting of a new codegen option behind-the-scenes, and small logic changes to the optimization pipeline to factor in the new flag. Existing tests all still pass.
- Loading branch information
1 parent
4713aeb
commit d164fd9
Showing
11 changed files
with
165 additions
and
91 deletions.
There are no files selected for viewing
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
43 changes: 43 additions & 0 deletions
43
clang/test/SemaSYCL/sycl-force-inline-kernel-lambda-ast.cpp
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,43 @@ | ||
| // RUN: %clang_cc1 -fno-sycl-force-inline-kernel-lambda -fsycl-is-device -internal-isystem %S/Inputs -disable-llvm-passes -triple spir64-unknown-unknown -ast-dump -o - %s | FileCheck %s --check-prefixes=NOINLINE,CHECK | ||
| // RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -disable-llvm-passes -triple spir64-unknown-unknown -ast-dump -o - %s | FileCheck %s --check-prefixes=INLINE,CHECK | ||
|
|
||
| // Tests that the appropriate inlining attributes are added to kernel lambda functions, | ||
| // with no inline attribute being added when -fno-sycl-force-inline-kernel-lambda is set | ||
| // and attribute not explicitly provided. | ||
|
|
||
| #include "sycl.hpp" | ||
|
|
||
| int main() { | ||
| sycl::queue q; | ||
|
|
||
| q.submit([&](sycl::handler &h) { | ||
| // CHECK: LambdaExpr{{.*}}sycl-force-inline-kernel-lambda-ast.cpp:17 | ||
| // INLINE: AlwaysInlineAttr | ||
| // NOINLINE-NOT: AlwaysInlineAttr | ||
| h.parallel_for<class KernelName>([] {}); | ||
| }); | ||
|
|
||
| q.submit([&](sycl::handler &h) { | ||
| // CHECK: LambdaExpr{{.*}}sycl-force-inline-kernel-lambda-ast.cpp:23 | ||
| // CHECK: AlwaysInlineAttr | ||
| h.parallel_for<class KernelNameInline>([]() __attribute__((always_inline)) {}); | ||
| }); | ||
|
|
||
| q.submit([&](sycl::handler &h) { | ||
| // CHECK: LambdaExpr{{.*}}sycl-force-inline-kernel-lambda-ast.cpp:30 | ||
| // CHECK: NoInlineAttr | ||
| // CHECK-NOT: AlwaysInlineAttr | ||
| h.parallel_for<class KernelNameNoInline>([]() __attribute__((noinline)) {}); | ||
| }); | ||
|
|
||
| /// The flag is ignored for ESIMD kernels | ||
| q.submit([&](sycl::handler &h) { | ||
| // CHECK: LambdaExpr{{.*}}sycl-force-inline-kernel-lambda-ast.cpp:39 | ||
| // CHECK: SYCLSimdAttr | ||
| // CHECK-NOT: AlwaysInlineAttr | ||
| // CHECK-NOT: NoInlineAttr | ||
| h.parallel_for<class KernelNameESIMD>([]() __attribute__((sycl_explicit_simd)) {}); | ||
| }); | ||
|
|
||
| return 0; | ||
| } |
This file was deleted.
Oops, something went wrong.
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,71 @@ | ||
| // RUN: %clangxx -fsycl -Xclang -fsycl-is-device -fsycl-device-only -Xclang -fno-sycl-early-optimizations -S -emit-llvm %s -o - | FileCheck %s | ||
|
|
||
| // Tests for IR of device_has(aspect, ...) attribute and | ||
| // !sycl_used_aspects metadata | ||
| #include <sycl/sycl.hpp> | ||
|
|
||
| using namespace sycl; | ||
| queue q; | ||
|
|
||
| // CHECK: define weak_odr dso_local spir_kernel void @{{.*}}kernel_name_1{{.*}} !sycl_declared_aspects ![[ASPECTS1:[0-9]+]] !srcloc ![[SRCLOC1:[0-9]+]] {{.*}} | ||
|
|
||
| // CHECK: define dso_local spir_func void @{{.*}}func1{{.*}} !sycl_declared_aspects ![[ASPECTS1]] !srcloc ![[SRCLOC2:[0-9]+]] | ||
| // CHECK-SAME: !sycl_used_aspects ![[ASPECTS1]] | ||
| [[sycl::device_has(sycl::aspect::cpu)]] void func1() {} | ||
|
|
||
| // CHECK: define dso_local spir_func void @{{.*}}func2{{.*}} !sycl_declared_aspects ![[ASPECTS2:[0-9]+]] !srcloc ![[SRCLOC3:[0-9]+]] | ||
| // CHECK-SAME: !sycl_used_aspects ![[ASPECTS2]] | ||
| [[sycl::device_has(sycl::aspect::fp16, sycl::aspect::gpu)]] void func2() {} | ||
|
|
||
| // CHECK: define dso_local spir_func void @{{.*}}func3{{.*}} !sycl_declared_aspects ![[EMPTYASPECTS:[0-9]+]] !srcloc ![[SRCLOC4:[0-9]+]] { | ||
| [[sycl::device_has()]] void func3() {} | ||
|
|
||
| // CHECK: define linkonce_odr dso_local spir_func void @{{.*}}func4{{.*}} !sycl_declared_aspects ![[ASPECTS3:[0-9]+]] !srcloc ![[SRCLOC5:[0-9]+]] | ||
| // CHECK-SAME: !sycl_used_aspects ![[ASPECTS3]] | ||
| template <sycl::aspect Aspect> [[sycl::device_has(Aspect)]] void func4() {} | ||
|
|
||
| // CHECK: define dso_local spir_func void @{{.*}}func5{{.*}} !sycl_declared_aspects ![[ASPECTS1]] !srcloc ![[SRCLOC6:[0-9]+]] | ||
| // CHECK-SAME: !sycl_used_aspects ![[ASPECTS1]] | ||
| [[sycl::device_has(sycl::aspect::cpu)]] void func5(); | ||
| void func5() {} | ||
|
|
||
| constexpr sycl::aspect getAspect() { return sycl::aspect::cpu; } | ||
| // CHECK: define dso_local spir_func void @{{.*}}func6{{.*}} !sycl_declared_aspects ![[ASPECTS1]] !srcloc ![[SRCLOC7:[0-9]+]] | ||
| // CHECK-SAME: !sycl_used_aspects ![[ASPECTS1]] | ||
| [[sycl::device_has(getAspect())]] void func6() {} | ||
|
|
||
| class KernelFunctor { | ||
| public: | ||
| [[sycl::device_has(sycl::aspect::cpu)]] void operator()() const { | ||
| func1(); | ||
| func2(); | ||
| func3(); | ||
| func4<sycl::aspect::host>(); | ||
| func5(); | ||
| func6(); | ||
| } | ||
| }; | ||
|
|
||
| void foo() { | ||
| q.submit([&](handler &h) { | ||
| KernelFunctor f1; | ||
| h.single_task<class kernel_name_1>(f1); | ||
| // CHECK: define weak_odr dso_local spir_kernel void @{{.*}}kernel_name_2{{.*}} !sycl_declared_aspects ![[ASPECTS4:[0-9]+]] !srcloc ![[SRCLOC8:[0-9]+]] {{.*}} | ||
| h.single_task<class kernel_name_2>( | ||
| []() [[sycl::device_has(sycl::aspect::gpu)]] {}); | ||
| }); | ||
| } | ||
|
|
||
| // CHECK: [[ASPECTS1]] = !{i32 1} | ||
| // CHECK: [[SRCLOC1]] = !{i32 {{[0-9]+}}} | ||
| // CHECK: [[EMPTYASPECTS]] = !{} | ||
| // CHECK: [[SRCLOC2]] = !{i32 {{[0-9]+}}} | ||
| // CHECK: [[ASPECTS2]] = !{i32 5, i32 2} | ||
| // CHECK: [[SRCLOC3]] = !{i32 {{[0-9]+}}} | ||
| // CHECK: [[SRCLOC4]] = !{i32 {{[0-9]+}}} | ||
| // CHECK: [[ASPECTS3]] = !{i32 0} | ||
| // CHECK: [[SRCLOC5]] = !{i32 {{[0-9]+}}} | ||
| // CHECK: [[SRCLOC6]] = !{i32 {{[0-9]+}}} | ||
| // CHECK: [[SRCLOC7]] = !{i32 {{[0-9]+}}} | ||
| // CHECK: [[ASPECTS4]] = !{i32 2} | ||
| // CHECK: [[SRCLOC8]] = !{i32 {{[0-9]+}}} |