-
Notifications
You must be signed in to change notification settings - Fork 738
Commit
This commit does not belong to any branch on this repository, and may belong to a fork outside of the repository.
[SYCL] Insert annotation in
annotated_ptr::get() (#12343)
When properties like alignment is specified in a `annotated_ptr` type,
certain operators (like `[]`, `+=`, `++`) are disabled. This results in
loop code to be written as follows:
```
annotated_ptr<int, decltype(properties{...alignment<8>...})> ann_ptr;
...
int *p = ann_ptr.get(); // ann_ptr cannot be used in the for loop directly
for (int i = 0; i < n; i++) {
p[i] = i;
}
```
When getting the underlying pointer, the annotation gets lost, so does
the the possible optimization on the for-loop brought by the
annotated_ptr properties.
This PR includes changes on spec, header and clang compiler:
1. In `annotated_ptr` spec, update the spec for the `get()` function
2. In the `annotated_ptr` header, update the `get()` function by
inserting `llvm.ptr.annotation`, so that on the target machines like
FPGA for which clang FE only performs O0 optimization, the annotation
inserted can be preserved for the corresponding backends to perform
platform-specific optimizations. For the example above, the `alignment`
information can help the FPGA compiler to build aligned loads/stores.
3. In the clang compiler, the pass `CompileTimePropertiesPass` used to
always drop `alignment` from the annotation string. This PR changes this
behavior to dropping `alignment` only when the compiler finds
load/store/MemIntrinsics in the users of `llvm.ptr.annotation` and
applies the alignment to these instructions.- Loading branch information
Showing
8 changed files
with
139 additions
and
52 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
48 changes: 48 additions & 0 deletions
48
sycl/test/extensions/annotated_ptr/annotation_insertion.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,48 @@ | ||
| // RUN: %clangxx -fsycl-device-only -fsycl-targets=spir64_fpga -S -emit-llvm %s -o - | FileCheck %s | ||
|
|
||
| // Tests that `@llvm.ptr.annotation` is inserted when calling | ||
| // `annotated_ptr::get()` | ||
|
|
||
| #include "sycl/sycl.hpp" | ||
| #include <sycl/ext/intel/fpga_extensions.hpp> | ||
|
|
||
| #include <iostream> | ||
|
|
||
| // clang-format on | ||
|
|
||
| using namespace sycl; | ||
| using namespace ext::oneapi::experimental; | ||
| using namespace ext::intel::experimental; | ||
|
|
||
| // CHECK: @[[AnnStr:.*]] = private unnamed_addr addrspace(1) constant [19 x i8] c"{5921:\220\22}{44:\228\22}\00" | ||
|
|
||
| using ann_ptr_t1 = | ||
| annotated_ptr<int, decltype(properties(buffer_location<0>, alignment<8>))>; | ||
|
|
||
| struct MyIP { | ||
| ann_ptr_t1 a; | ||
|
|
||
| MyIP(int *a_) : a(a_) {} | ||
|
|
||
| void operator()() const { | ||
| // CHECK: %ptr.addr = alloca ptr addrspace(4), align 8 | ||
| // CHECK: store ptr addrspace(4) %ptr, ptr %ptr.addr, align 8 | ||
| // CHECK: %[[LoadPtr:.*]] = load ptr addrspace(4), ptr %ptr.addr, align 8 | ||
| // CHECK: %[[AnnPtr:.*]] = call ptr addrspace(4) @llvm.ptr.annotation.p4.p1(ptr addrspace(4) %[[LoadPtr]], ptr addrspace(1) @[[AnnStr]] | ||
| // CHECK: ret ptr addrspace(4) %[[AnnPtr]] | ||
| int *ptr = a.get(); // llvm.ptr.annotation is inserted | ||
| *ptr = 15; | ||
| } | ||
| }; | ||
|
|
||
| void TestVectorAddWithAnnotatedMMHosts() { | ||
| sycl::queue q; | ||
| auto raw = malloc_shared<int>(5, q); | ||
| q.submit([&](handler &h) { h.single_task(MyIP{raw}); }).wait(); | ||
| free(raw, q); | ||
| } | ||
|
|
||
| int main() { | ||
| TestVectorAddWithAnnotatedMMHosts(); | ||
| return 0; | ||
| } |
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