Skip to content
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

Added new IfConditionalConversion to fold nested and-also and or-else block chains #328

Merged
merged 11 commits into from Dec 4, 2020

Conversation

m4rs-mt
Copy link
Owner

@m4rs-mt m4rs-mt commented Nov 22, 2020

Currently, ILGPU cannot fold nested && and || CFG-based block jumps into logical & and | operations. This PR adds a new transformation called IfConditionalConversion, which converts such block jumps into efficient logical operations without branches.
Consider the sample below:

if (a > 0 && b < c || d > 4) {
   ...
}

which can be compiled using v0.9.2 + O2 into the following sequence of PTX instructions:

L_9710:
setp.le.s32	%p1, %r1, 0;
@%p1 bra	L_9712;
bra	L_9711;

L_9711:
setp.lt.s32	%p2, %r2, %r3;
@%p2 bra	L_9713;
bra	L_9712;

L_9712:
setp.le.s32	%p3, %r4, 4;
@%p3 bra	L_9714;
bra	L_9713;

L_9713:
mov.b32	%r5, %ctaid.x;
mov.b32	%r6, %tid.x;
mov.b32	%r7, %ntid.x;
mul.lo.s32	%r8, %r5, %r7;
add.s32	%r9, %r6, %r8;
mul.wide.u32	%rd4, %r9, 4;
add.u64	%rd3, %rd1, %rd4;
st.b32	[%rd3], 42;
bra	L_9714;

With the help of this PR this piece of code is then converted into (since all blocks are considered "small enough" and have no side effects):

if (a > 0 & b < c | d > 4) {
   ...
}

which results in the following sequence of PTX instructions:

setp.gt.s32	%p1, %r1, 0;
setp.lt.s32	%p2, %r2, %r3;
and.pred	%p3, %p1, %p2;
setp.gt.s32	%p4, %r4, 4;
or.pred	%p5, %p3, %p4;
@%p5 bra	L_11021;

L_11020:
...

L_11021:
mov.b32	%r5, %ctaid.x;
mov.b32	%r6, %tid.x;
mov.b32	%r7, %ntid.x;
mul.lo.s32	%r8, %r5, %r7;
add.s32	%r9, %r6, %r8;
mul.wide.u32	%rd4, %r9, 4;
add.u64	%rd3, %rd1, %rd4;
st.global.b32	[%rd3], 42;
bra	L_11020;

@m4rs-mt m4rs-mt added this to the v0.9.3 milestone Nov 22, 2020
@m4rs-mt m4rs-mt marked this pull request as draft November 22, 2020 13:44
@m4rs-mt m4rs-mt marked this pull request as ready for review November 26, 2020 18:43
Src/ILGPU/IR/BasicBlockMapping.cs Show resolved Hide resolved
@m4rs-mt m4rs-mt merged commit 3101458 into master Dec 4, 2020
@m4rs-mt m4rs-mt deleted the if_conditionals branch December 4, 2020 12:27
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

None yet

2 participants