Repository navigation
[SYCL] Implement basic sub-buffers support - #64
Conversation
Signed-off-by: Mariya Podchishchaeva <[email protected]>
keryell
left a comment
There was a problem hiding this comment.
Nice to have this new feature! \o/
Just a few things that could be clarified.
| const int N = 7; | ||
| int Result[M][N] = {0}; | ||
| { | ||
| auto OrigRange = range<2>(M, N); |
There was a problem hiding this comment.
What about using modern C++ { } everywhere instead of old troublesome ()?
Specially if people look at tests as good coding example and if we consider that SYCL is about modern C++ and heterogeneous computing... :-)
| { | ||
| auto OrigRange = range<2>(M, N); | ||
| buffer<int, 2> Buffer(OrigRange); | ||
| Buffer.set_final_data((int *)Result); |
There was a problem hiding this comment.
Old C cast in modern C++ code...
| bool Failed = false; | ||
| // Basic test case | ||
| { | ||
| const int M = 6; |
There was a problem hiding this comment.
I think you can replace all the const something in this code by constexpr something
| for (size_t i = 0; i < M; ++i) { | ||
| for (size_t j = 0; j < N; ++j) { | ||
| size_t Expected = | ||
| ((i == 0) || (i == M - 1) || (j == 0) || (j == N - 1)) ? 0 : 1; |
There was a problem hiding this comment.
Not clear why a size_t. Actually a bool looks fine to me.
auto Expected = !(i == 0) || (i == M - 1) || (j == 0) || (j == N - 1));
Replacing auto by int is also possible if you think it would cause the reader more time to understand the next line...
| { | ||
| const int M = 10; | ||
| int Data[M] = {0}; | ||
| auto OrigRange = range<1>(M); |
There was a problem hiding this comment.
Shorter:
range<1> OrigRange { M };
| int Data[M] = {0}; | ||
| auto OrigRange = range<1>(M); | ||
| buffer<int, 1> Buffer(Data, OrigRange); | ||
| auto Offset = id<1>(1); |
| }); | ||
| auto Acc = Buffer.get_access<cl::sycl::access::mode::read>(); | ||
| for (size_t i = 0; i < M; ++i) { | ||
| size_t Expected = (i > 1 && i < M - 2) ? 1 : 0; |
There was a problem hiding this comment.
int Expected = i > 1 && i < M - 2;
| buffer(const range<dimensions> &bufferRange, | ||
| const property_list &propList = {}) | ||
| : Range(bufferRange) { | ||
| : Range(bufferRange), MemRange(bufferRange) { |
There was a problem hiding this comment.
What about writing modern C++ with some { } instead of () for most of the instance constructions?
Summary: According to the new Armv8-M specification https://static.docs.arm.com/ddi0553/bh/DDI0553B_h_armv8m_arm.pdf the instructions SQRSHRL and UQRSHLL now have an additional immediate operand <saturate>. The new assembly syntax is: SQRSHRL<c> RdaLo, RdaHi, #<saturate>, Rm UQRSHLL<c> RdaLo, RdaHi, #<saturate>, Rm where <saturate> can be either 64 (the existing behavior) or 48, in that case the result is saturated to 48 bits. The new operand is encoded as follows: #64 Encoded as sat = 0 #48 Encoded as sat = 1 sat is bit 7 of the instruction bit pattern. This patch adds a new assembler operand class MveSaturateOperand which implements parsing and encoding. Decoding is implemented in DecodeMVEOverlappingLongShift. Reviewers: ostannard, simon_tatham, t.p.northover, samparker, dmgreen, SjoerdMeijer Reviewed By: simon_tatham Subscribers: javed.absar, kristof.beyls, hiraditya, pbarrio, llvm-commits Tags: #llvm Differential Revision: https://reviews.llvm.org/D64810 llvm-svn: 366555
Summary: feature coverage is a useful signal that is available during the merge process, but was not printed previously. Output example: ``` $ ./fuzzer -use_value_profile=1 -merge=1 new_corpus/ seed_corpus/ INFO: Seed: 1676551929 INFO: Loaded 1 modules (2380 inline 8-bit counters): 2380 [0x90d180, 0x90dacc), INFO: Loaded 1 PC tables (2380 PCs): 2380 [0x684018,0x68d4d8), MERGE-OUTER: 180 files, 78 in the initial corpus MERGE-OUTER: attempt 1 INFO: Seed: 1676574577 INFO: Loaded 1 modules (2380 inline 8-bit counters): 2380 [0x90d180, 0x90dacc), INFO: Loaded 1 PC tables (2380 PCs): 2380 [0x684018,0x68d4d8), INFO: -max_len is not provided; libFuzzer will not generate inputs larger than 1048576 bytes MERGE-INNER: using the control file '/tmp/libFuzzerTemp.111754.txt' MERGE-INNER: 180 total files; 0 processed earlier; will process 180 files now #1 pulse cov: 134 ft: 330 exec/s: 0 rss: 37Mb #2 pulse cov: 142 ft: 462 exec/s: 0 rss: 38Mb #4 pulse cov: 152 ft: 651 exec/s: 0 rss: 38Mb #8 pulse cov: 152 ft: 943 exec/s: 0 rss: 38Mb #16 pulse cov: 520 ft: 2783 exec/s: 0 rss: 39Mb #32 pulse cov: 552 ft: 3280 exec/s: 0 rss: 41Mb #64 pulse cov: 576 ft: 3641 exec/s: 0 rss: 50Mb #78 LOADED cov: 602 ft: 3936 exec/s: 0 rss: 88Mb #128 pulse cov: 611 ft: 3996 exec/s: 0 rss: 93Mb #180 DONE cov: 611 ft: 4016 exec/s: 0 rss: 155Mb MERGE-OUTER: succesfull in 1 attempt(s) MERGE-OUTER: the control file has 39741 bytes MERGE-OUTER: consumed 0Mb (37Mb rss) to parse the control file MERGE-OUTER: 9 new files with 80 new features added; 9 new coverage edges ``` Reviewers: hctim, morehouse Reviewed By: morehouse Subscribers: delcypher, #sanitizers, llvm-commits, kcc Tags: #llvm, #sanitizers Differential Revision: https://reviews.llvm.org/D66030 llvm-svn: 368617
Summary:
The greedy register allocator occasionally decides to insert a large number of
unnecessary copies, see below for an example. The -consider-local-interval-cost
option (which X86 already enables by default) fixes this. We enable this option
for AArch64 only after receiving feedback that this change is not beneficial for
PowerPC.
We evaluated the impact of this change on compile time, code size and
performance benchmarks.
This option has a small impact on compile time, measured on CTMark. A 0.1%
geomean regression on -O1 and -O2, and 0.2% geomean for -O3, with at most 0.5%
on individual benchmarks.
The effect on both code size and performance on AArch64 for the LLVM test suite
is nil on the geomean with individual outliers (ignoring short exec_times)
between:
best worst
size..text -3.3% +0.0%
exec_time -5.8% +2.3%
On SPEC CPU® 2017 (compiled for AArch64) there is a minor reduction (-0.2% at
most) in code size on some benchmarks, with a tiny movement (-0.01%) on the
geomean. Neither intrate nor fprate show any change in performance.
This patch makes the following changes.
- For the AArch64 target, enableAdvancedRASplitCost() now returns true.
- Ensures that -consider-local-interval-cost=false can disable the new
behaviour if necessary.
This matrix multiply example:
$ cat test.c
long A[8][8];
long B[8][8];
long C[8][8];
void run_test() {
for (int k = 0; k < 8; k++) {
for (int i = 0; i < 8; i++) {
for (int j = 0; j < 8; j++) {
C[i][j] += A[i][k] * B[k][j];
}
}
}
}
results in the following generated code on AArch64:
$ clang --target=aarch64-arm-none-eabi -O3 -S test.c -o -
[...]
// %for.cond1.preheader
// =>This Inner Loop Header: Depth=1
add x14, x11, x9
str q0, [sp, #16] // 16-byte Folded Spill
ldr q0, [x14]
mov v2.16b, v15.16b
mov v15.16b, v14.16b
mov v14.16b, v13.16b
mov v13.16b, v12.16b
mov v12.16b, v11.16b
mov v11.16b, v10.16b
mov v10.16b, v9.16b
mov v9.16b, v8.16b
mov v8.16b, v31.16b
mov v31.16b, v30.16b
mov v30.16b, v29.16b
mov v29.16b, v28.16b
mov v28.16b, v27.16b
mov v27.16b, v26.16b
mov v26.16b, v25.16b
mov v25.16b, v24.16b
mov v24.16b, v23.16b
mov v23.16b, v22.16b
mov v22.16b, v21.16b
mov v21.16b, v20.16b
mov v20.16b, v19.16b
mov v19.16b, v18.16b
mov v18.16b, v17.16b
mov v17.16b, v16.16b
mov v16.16b, v7.16b
mov v7.16b, v6.16b
mov v6.16b, v5.16b
mov v5.16b, v4.16b
mov v4.16b, v3.16b
mov v3.16b, v1.16b
mov x12, v0.d[1]
fmov x15, d0
ldp q1, q0, [x14, #16]
ldur x1, [x10, #-256]
ldur x2, [x10, #-192]
add x9, x9, #64 // =64
mov x13, v1.d[1]
fmov x16, d1
ldr q1, [x14, #48]
mul x3, x15, x1
mov x14, v0.d[1]
fmov x17, d0
mov x18, v1.d[1]
fmov x0, d1
mov v1.16b, v3.16b
mov v3.16b, v4.16b
mov v4.16b, v5.16b
mov v5.16b, v6.16b
mov v6.16b, v7.16b
mov v7.16b, v16.16b
mov v16.16b, v17.16b
mov v17.16b, v18.16b
mov v18.16b, v19.16b
mov v19.16b, v20.16b
mov v20.16b, v21.16b
mov v21.16b, v22.16b
mov v22.16b, v23.16b
mov v23.16b, v24.16b
mov v24.16b, v25.16b
mov v25.16b, v26.16b
mov v26.16b, v27.16b
mov v27.16b, v28.16b
mov v28.16b, v29.16b
mov v29.16b, v30.16b
mov v30.16b, v31.16b
mov v31.16b, v8.16b
mov v8.16b, v9.16b
mov v9.16b, v10.16b
mov v10.16b, v11.16b
mov v11.16b, v12.16b
mov v12.16b, v13.16b
mov v13.16b, v14.16b
mov v14.16b, v15.16b
mov v15.16b, v2.16b
ldr q2, [sp] // 16-byte Folded Reload
fmov d0, x3
mul x3, x12, x1
[...]
With -consider-local-interval-cost the same section of code results in the
following:
$ clang --target=aarch64-arm-none-eabi -mllvm -consider-local-interval-cost -O3 -S test.c -o -
[...]
.LBB0_1: // %for.cond1.preheader
// =>This Inner Loop Header: Depth=1
add x14, x11, x9
ldp q0, q1, [x14]
ldur x1, [x10, #-256]
ldur x2, [x10, #-192]
add x9, x9, #64 // =64
mov x12, v0.d[1]
fmov x15, d0
mov x13, v1.d[1]
fmov x16, d1
ldp q0, q1, [x14, #32]
mul x3, x15, x1
cmp x9, #512 // =512
mov x14, v0.d[1]
fmov x17, d0
fmov d0, x3
mul x3, x12, x1
[...]
Reviewers: SjoerdMeijer, samparker, dmgreen, qcolombet
Reviewed By: dmgreen
Subscribers: ZhangKang, jsji, wuzish, ppc-slack, lkail, steven.zhang, MatzeB, qcolombet, kristof.beyls, hiraditya, llvm-commits
Tags: #llvm
Differential Revision: https://reviews.llvm.org/D69437
[compiler] Add a pass to replace target extension types
Swift async functions can miscompile on Windows ARM64 at `-O2`, when there's enough register pressure that a local gets scavenged into the callee-save area: the local ends up sharing an address with the saved caller x29, so the epilogue restores a value the function has already overwritten. swiftlang/swift#90920 has a reduced repro. `assignCalleeSavedSpillSlots` creates the swift async context object before the callee-save loop instead of inside it next to the FP slot. MachineFrameInfo ends up with it above the frame record while the prologue stores it below at FP-8, and the 8 byte disagreement leaves a hole in the middle of the callee-save area. PEI's scavenger hands that hole to the local. Only reproduces at -O2 and up since scavenging is gated on the opt level. ``` sub sp, sp, #112 str x19, [sp, #16] // 8-byte Spill str x21, [sp, #24] // 8-byte Spill stp x23, x24, [sp, #32] // 16-byte Folded Spill stp x25, x26, [sp, #48] // 16-byte Folded Spill stp x27, x28, [sp, #64] // 16-byte Folded Spill stp x29, x30, [sp, #88] // 16-byte Folded Spill str xzr, [sp, #80] add x29, sp, #88 ... str x7, [x29] // 8-byte Spill ... ldr x1, [x29] // 8-byte Reload ldp x29, x30, [sp, #88] // 16-byte Folded Reload ``` This creates the object inside the loop so the two agree. The other option was leaving the creation site alone and teaching MachineFrameInfo about the expanded 24 byte FP/LR slot, but that puts the layout in two places. Not sure which is preferred here, I don't know this code well. This also asserts the saved FP object resolves to FP+0, since nothing checks that today. Reverting the fix makes it fire on the same funclet. `store-swift-async-context-clobber-live-reg.ll` already miscompiles with `-regalloc=fast`, so this isn't Swift specific.
`BigInt::operator*` always computes the full 2N-word product (`ful_mul`, i.e. `multiword::multiply_with_carry`) and then truncates. In order to scalarize the process to generate efficient code, the loop needs to be fully unrolled, which fails on 32-bit targets for int128 due to LLVM's default threshold. This patch adjust the hint to allow int128 to be unrolled while gating larger integers away. ### 64-bit targets: roughly okay aarch64 is byte-identical before/after; x86-64 gets one `mov` shorter. After the patch, 128×128→128 (lo/hi in memory): ```asm ; aarch64 ; x86-64 ldp x8, x12, [x0] movq (%rdi), %r9 ldp x9, x10, [x1] movq (%rsi), %r8 mul x10, x10, x8 movq %r8, %rax mul x11, x9, x8 mulq %r9 umulh x8, x9, x8 imulq 8(%rsi), %r9 madd x9, x9, x12, x10 imulq 8(%rdi), %r8 add x8, x9, x8 addq %r9, %r8 stp x11, x8, [x2] addq %rdx, %r8 ret ... ``` ### 32-bit targets (e.g. cortex-m85): badly bloated A 64-byte stack frame, carry selection through IT blocks, and a call into an outlined `ful_mul`. Kernel is the folded multiply (`lo ^ hi` of the `UInt<128>` product): ```asm ; before ; after sub sp, #64 umull r12, lr, r2, r0 mov.w r12, #0 umull r2, r4, r2, r1 strd r12, r12, [sp, #24] umull r0, r5, r3, r0 orrs.w lr, r0, r1 umaal r4, r5, r3, r1 it ne adds.w r1, lr, r2 strdne r0, r1, [sp, #16] adcs r2, r4, #0 ... (5 more stores) adc r3, r5, #0 bl BigInt<128>::ful_mul adds r0, r0, r1 ldm r3, {r0, r1, r2, r3} adcs r2, r2, #0 eors r1, r3 adc r1, r3, #0 eors r0, r2 eors r1, r0 add sp, #64 eor.w r0, r2, r12 ``` llvm-mca (cortex-m85): **114 → 11 cycles** per folded multiply. ### Benchmarks | environment | before | after | speedup | |---|---|---|---| | EK-RA8M2 (Cortex-M85 r1p1, DWT cycle counter, I-cache on) | 250 cyc/mul | 20 cyc/mul | 12.5× | | wasm32 under wasmer 7.3, `-U__SIZEOF_INT128__` (see notes) | 0.70 s / 100M muls | 0.28 s | 2.5× | Bench kernel (dependent chain, identical results before/after): ```cpp using B = LIBC_NAMESPACE::BigInt<128, false, uint64_t>; B x = (B(a) << 64) | B(~a); B y = (B(b) << 64) | B(~b); for (int i = 0; i < n; i++) { x = x * y; y.val[0] ^= x.val[1]; } return x.val[0] ^ x.val[1]; ``` ### Carry builtin dispatch fix Additionally, fix a typo in the `add_with_carry` / `sub_with_borrow` builtin dispatch: the `#if`/`#elif` chain over `__has_builtin`. ### Notes 1. Another workaround for such targets is to force-enable native 128-bit integers (clang's `-fforce-enable-int128`) 2. The wasm number above is benched with 128-bit integers forcibly disabled (`-U__SIZEOF_INT128__`), just to check its effect. --- Guided and reviewed by human Co-Authored-By: Claude Fable 5 <[email protected]> --------- Co-authored-by: Claude Fable 5 <[email protected]>
Swift async functions can miscompile on Windows ARM64 at `-O2`, when there's enough register pressure that a local gets scavenged into the callee-save area: the local ends up sharing an address with the saved caller x29, so the epilogue restores a value the function has already overwritten. swiftlang/swift#90920 has a reduced repro. `assignCalleeSavedSpillSlots` creates the swift async context object before the callee-save loop instead of inside it next to the FP slot. MachineFrameInfo ends up with it above the frame record while the prologue stores it below at FP-8, and the 8 byte disagreement leaves a hole in the middle of the callee-save area. PEI's scavenger hands that hole to the local. Only reproduces at -O2 and up since scavenging is gated on the opt level. ``` sub sp, sp, #112 str x19, [sp, #16] // 8-byte Spill str x21, [sp, #24] // 8-byte Spill stp x23, x24, [sp, #32] // 16-byte Folded Spill stp x25, x26, [sp, #48] // 16-byte Folded Spill stp x27, x28, [sp, #64] // 16-byte Folded Spill stp x29, x30, [sp, #88] // 16-byte Folded Spill str xzr, [sp, #80] add x29, sp, #88 ... str x7, [x29] // 8-byte Spill ... ldr x1, [x29] // 8-byte Reload ldp x29, x30, [sp, #88] // 16-byte Folded Reload ``` This creates the object inside the loop so the two agree. The other option was leaving the creation site alone and teaching MachineFrameInfo about the expanded 24 byte FP/LR slot, but that puts the layout in two places. Not sure which is preferred here, I don't know this code well. This also asserts the saved FP object resolves to FP+0, since nothing checks that today. Reverting the fix makes it fire on the same funclet. `store-swift-async-context-clobber-live-reg.ll` already miscompiles with `-regalloc=fast`, so this isn't Swift specific. (cherry picked from commit c358e8d)
Signed-off-by: Mariya Podchishchaeva [email protected]