# Problem with left shift

**URL:** <https://forums.developer.nvidia.com/t/problem-with-left-shift/20660>\
**Category:** CUDA Programming and Performance\
**Created:** [January 21, 2011, 5:12pm UTC](https://forums.developer.nvidia.com/t/problem-with-left-shift/20660 "2011-01-21T17:12:39Z")\
**Posts on this page:** 5\
**Page:** 1

<div class="post-metadata">

**Author:** ![GiuseppeC](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@GiuseppeC](https://forums.developer.nvidia.com/u/GiuseppeC)\
**Post date:** [January 21, 2011, 5:12pm UTC](https://forums.developer.nvidia.com/t/problem-with-left-shift/20660/1 "2011-01-21T17:12:39Z")

</div>

I’ve a problem with the “\<\<” operator (left shift).

Somewhere in my code there is the following define:

```auto
#define ROL64(a, offset) ((offset != 0) ? ((((UINT64)a) << offset) ^ (((UINT64)a) >> (64-offset))) : a)

```

When I try to compile my program I obtain a very strange error:

```auto
### Assertion failure at line 1923 of ../../be/cg/cgemit.cxx:

### Compiler Error in file /tmp/tmpxft_000022e5_00000000-7_kernel.cpp3.i during Assembly phase:

### incorrect register class for operand 1

nvopencc INTERNAL ERROR: /opt/cuda/open64/lib//be returned non-zero status 1

```

This error disappears if, instead of using the “\<\<”, I use the “\>\>” (right shift). I do not understand this behavior. By now I’m using the following “trick”

```auto
#define ROL64(a, offset) ((offset != 0) ? ((((UINT64)a) >> (0-offset)) ^ (((UINT64)a) >> (64-offset))) : a)

```

but I’m not sure this has the expected behavior.

Any Ideas?

---

<div class="post-metadata">

**Author:** ![SPWorley](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@SPWorley](https://forums.developer.nvidia.com/u/SPWorley)\
**Post date:** [January 21, 2011, 6:23pm UTC](https://forums.developer.nvidia.com/t/problem-with-left-shift/20660/2 "2011-01-21T18:23:45Z")

</div>

I submitted a similar bug to NVIDIA a few weeks ago which is clearly a compiler bug. This is probably the same one. It had to do with a shifts of 64 bit long values when the shift had integer math in it.

It only failed for sm\_20 and not sm\_13.

As a workaround, just put the shift amount into a temporary variable.

```auto
// compile with nvcc -arch sm_20 crash.cu

__device__ unsigned long circularShift(unsigned long v, int shift)

{

    return (v<<(64-shift));

}

int main() {}

```

---

<div class="post-metadata">

**Author:** ![njuffa](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@njuffa](https://forums.developer.nvidia.com/u/njuffa)\
**Post date:** [January 21, 2011, 7:12pm UTC](https://forums.developer.nvidia.com/t/problem-with-left-shift/20660/3 "2011-01-21T19:12:19Z")

</div>

I recall a compiler bug being reported regarding 64-bit rotates recently. If memory serves, one workaround is to assign the left-shift and right-shift portions to temporary variables, then combine them (with | or + for example) in a separate step. Unfortunately that won’t work for a macro due to the need for temp variables. Could you use a **forceinline** function instead?

BTW, negative shift counts cause undefined behavior in C/C++, so assuming that offset is \>= 0, ((a) \>\> (0-offset)) will not in general give the desired result.

---

<div class="post-metadata">

**Author:** ![SPWorley](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@SPWorley](https://forums.developer.nvidia.com/u/SPWorley)\
**Post date:** [January 21, 2011, 10:15pm UTC](https://forums.developer.nvidia.com/t/problem-with-left-shift/20660/4 "2011-01-21T22:15:47Z")

</div>

Just as an aside, since shifts aren’t a common topic usually worth posting for: GPU shifts behave differently than CPU shifts when the shift amount exceeds the word size in bits.

This is undefined behavior in the C standard, and is one of the few cases where NVIDIA GPUs differ in their integer computation from Intel and AMD CPUs.

On the CPUs, a shift amount \>= word size will end up using the low order bits of the shift, so a integer shift of 33 is the same as a shift of 1, for example.  
On the GPU, the shift value is saturated, so a shift of 32 or more acts like 32.  
This is true for right or left shifts.

Again this is not a bug, it’s documented as undefined behavior in C itself. It’s interesting there’s a difference though.

---

<div class="post-metadata">

**Author:** ![GiuseppeC](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@GiuseppeC](https://forums.developer.nvidia.com/u/GiuseppeC)\
**Post date:** [January 24, 2011, 5:56pm UTC](https://forums.developer.nvidia.com/t/problem-with-left-shift/20660/5 "2011-01-24T17:56:22Z")

</div>

Thank you very much for the detailed informations :)

I modified my code in this way:

```auto
// NVCC Bug

//#define ROL64(a, offset) ((offset != 0) ? ((((UINT64)a) << offset) ^ (((UINT64)a) >> (64-offset))) : a)

__device__ inline UINT64 ROL64(UINT64 a, unsigned int offset)

{

        const int _offset = offset;

        return ((offset != 0) ? ((a << _offset) ^ (a >> (64-offset))) : a);

}

```

By now I hope it will be ok, however I have to check the range of the variable “offset” in order to avoid misalignment between the CPU and GPU code.