[WIP] Add bitwise operators and propagate them from parser -> halide -> isl

[prigoyal]

trying to add bitwise operators. so far they are parsed correctly in their precedence order. tc2halide conversion is also done.

I am not sure about how to propagate these through to isl yet.

NOTE: I stacked this PR on add-operators branch to avoid merge conflicts.

closes #131

[prigoyal]

@skimo-openhub , @ftynse , how do the operators propagate to isl? I see files like tc/external/detail/islpp.h but I don't fully understand how/when/if operators are added there. Can you provide some guidance? Thanks a lot.

[abadams]

Note that the Halide simplifier will convert all the simpler bit operations to multiplication (x << 5), division (x >> 5), and modulo (x & 0x3f), so the only bit operations that reach isl will be the nastier ones. That said, bit operations do have well-defined bounds that we may be able to exploit (e.g. for unsigned ints (x & y) <= x)

[prigoyal]

thanks @abadams . I have two questions for @skimo-openhub and @ftynse:

1. should we support quasi-affine bitwise or can it be more general? from looking at halide bitwise ops, https://github.com/halide/Halide/blob/master/test/correctness/bitwise_ops.cpp, it looks like it doesn't need to be quasi-affine but not sure about isl
2. any ideas around how these ops will be propagated to isl? do we need to make an entry in tc/external/detail/islpp.h?

I'll try to add simple unit tests for these to catch errors on propagation to isl

[ftynse]

(quick answer, I'll give more details when not on the phone)

We need to differentiate two cases for bitwise operations, similarly to modulo. If they appear in tensor subscripts, those are not affine and should be overapproximated. Outer operations should be handled as function calls essentially.

[ftynse]

random thought: I would add checks in Sema that bitwise operations are applied to integer types, I don't think we want them defined for floats

[prigoyal]

thanks. I will wait to get more details to understand better the cases to distinguish.

In the meantime, propagation from halide to isl doesn't work by default. here is the codegen for one example:

extern "C" {
__global__ void f(int32* pc, const int32* pa, const int32* pb) {
  int b0 = blockIdx.x; int b1 = blockIdx.y; int b2 = blockIdx.z;
  int t0 = threadIdx.x; int t1 = threadIdx.y; int t2 = threadIdx.z;
  int32 (*c) = reinterpret_cast<int32 (*)>(pc);
  const int32 (*a) = reinterpret_cast<const int32 (*)>(pa);
  const int32 (*b) = reinterpret_cast<const int32 (*)>(pb);
  c[0] = shift_left(a[0], b[0]);
}
}

Note the shift_left is the halide representation but I don't know what these operations are called in isl. I would appreciate any pointers. cc @skimo-openhub

[ftynse]

isl must be agnostic to these operations. Essentially, the level we are interested about in isl is below those. At that level, there is no difference between A[i] + B[i] and A[i] << B[i]. We are only interested by what data is accessed, not which operations are performed on it. (This get different when expressions appear in the subscripts because it affects what data is accessed, but let's punt on it for now.)

In this particular case, the simplest solution is to inject #define shift_left(a,b) ((a) << (b)) in the code to make Halide's function work. I don't like this solution because it makes shift_left a reserved name in TC, but we already do so for something else if I'm not mistaken. Alternatively, somewhere on isl/Halide codegen boundary, we could rewrite that expression to use <<. Unfortunately, I cannot give you a pointer there immediately.

[ftynse]

Now, if operations appear in subscripts, like O(i << j) = A(i), isl must care about those because it changes what data in O is accessed. isl in particular and the polyhedral model in general is mostly designed for handling affine expressions (that is, sums of constants and products of constants with variables). There are some extensions allowed, which we call quasi-affine expressions. More specifically, division and modulus by a constant is allowed. Whenever the subscript expression is not affine, the common approach in the polyhedral is to over-approximate it: instead of saying "i-th iteration accesses O[i]", we say "i-th iteration may access any element of O".

Bitiwise operations in general are not affine, so you would need that over-approximation. It should be rather straightforward to implement in makeIslAffFrom... functions, there are already some non-affine cases. Shifts with constant RHS are a special case. They can be transformed to products (left shift) and divisions (right shift) with a constant factor, and that is (quasi)affine.

[prigoyal]

thanks @ftynse, the bitwise operators work now for the simple test case added to check correctness.

can you provide what more test cases should I test this on (especially over-approximation, affine/non-affine scenarios)? I believe that would help me understand more how to address those cases by digging more into the codebase :) thanks a lot for your guidance.

[skimo-openhub]

add me back when there is any progress

[facebook-github-bot]

Thank you for your pull request. We require contributors to sign our Contributor License Agreement, and yours has expired.

Before we can review or merge your code, we need you to email [email protected] with your details so we can update your status.

[facebook-github-bot]

Thank you for signing our Contributor License Agreement. We can now accept your code for this (and any) Facebook open source project. Thanks!

[facebook-github-bot]

Thank you for signing our Contributor License Agreement. We can now accept your code for this (and any) Facebook open source project. Thanks!