# Two questions (I'm a newbie)

**URL:** <https://forums.developer.nvidia.com/t/two-questions-im-a-newbie/8545>\
**Category:** CUDA Programming and Performance\
**Created:** [March 25, 2009, 5:14pm UTC](https://forums.developer.nvidia.com/t/two-questions-im-a-newbie/8545 "2009-03-25T17:14:24Z")\
**Posts on this page:** 10\
**Page:** 1

<div class="post-metadata">

**Author:** ![Noel\_Lopes](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@Noel\_Lopes](https://forums.developer.nvidia.com/u/Noel_Lopes)\
**Post date:** [March 25, 2009, 5:14pm UTC](https://forums.developer.nvidia.com/t/two-questions-im-a-newbie/8545/1 "2009-03-25T17:14:24Z")

</div>

Since the modulo operator is very expensive on GPUs:

I’m wondering if replacing this:

if (threadIdx.x % nextInterval == 0) // …

by this (I know that this will work only in a few cases, but that’s not my point) :

if ((threadIdx.x & (nextInterval-1)) == 0) // …

is faster?

Also I was wondering if:

x \<\<= 1;

is faster than:

x \*= 2;

---

<div class="post-metadata">

**Author:** ![pvouzis](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@pvouzis](https://forums.developer.nvidia.com/u/pvouzis)\
**Post date:** [March 25, 2009, 5:20pm UTC](https://forums.developer.nvidia.com/t/two-questions-im-a-newbie/8545/2 "2009-03-25T17:20:03Z")

</div>

Of course, both of them are faster, as they would be in any processor.

---

<div class="post-metadata">

**Author:** ![Paul\_Russell](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@Paul\_Russell](https://forums.developer.nvidia.com/u/Paul_Russell)\
**Post date:** [March 26, 2009, 1:14pm UTC](https://forums.developer.nvidia.com/t/two-questions-im-a-newbie/8545/3 "2009-03-26T13:14:09Z")

</div>

The first example will only work if nextInterval is a power of 2, of course, but otherwise is a useful optimisation.

For the second example any decent compiler will generate identical code for i\*2 and i\<\<1.

---

<div class="post-metadata">

**Author:** ![seibert](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/seibert/32/520688_2.png) [@seibert](https://forums.developer.nvidia.com/u/seibert)\
**Post date:** [March 26, 2009, 2:12pm UTC](https://forums.developer.nvidia.com/t/two-questions-im-a-newbie/8545/4 "2009-03-26T14:12:38Z")

</div>

> [@](#):
>
> For the second example any decent compiler will generate identical code for i\*2 and i\<\<1.

I checked this briefly with nvcc 2.1, and at the PTX level, this optimization is not performed. I don’t know if ptxas performs this optimization because I don’t have decuda handy.

---

<div class="post-metadata">

**Author:** ![Paul\_Russell](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@Paul\_Russell](https://forums.developer.nvidia.com/u/Paul_Russell)\
**Post date:** [March 27, 2009, 1:47pm UTC](https://forums.developer.nvidia.com/t/two-questions-im-a-newbie/8545/5 "2009-03-27T13:47:30Z")

</div>

> [@](#):
>
> I checked this briefly with nvcc 2.1, and at the PTX level, this optimization is not performed. I don’t know if ptxas performs this optimization because I don’t have decuda handy.

Perhaps if the cost of a multiply and a shift is the same then there is no need for the optimization. One tends to assume that a multiply will be more expensive on a conventional CPU but this may not be true on a GPU core.

---

<div class="post-metadata">

**Author:** ![seibert](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/seibert/32/520688_2.png) [@seibert](https://forums.developer.nvidia.com/u/seibert)\
**Post date:** [March 27, 2009, 2:35pm UTC](https://forums.developer.nvidia.com/t/two-questions-im-a-newbie/8545/6 "2009-03-27T14:35:52Z")

</div>

It is using the mul.lo.s32 instruction, which according to the programming guide, is actually 4 times slower than a bitwise operation. If ptxas isn’t performing this optimization, this is a pretty obvious thing that nvcc should be doing.

---

<div class="post-metadata">

**Author:** ![Paul\_Russell](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@Paul\_Russell](https://forums.developer.nvidia.com/u/Paul_Russell)\
**Post date:** [March 30, 2009, 12:29pm UTC](https://forums.developer.nvidia.com/t/two-questions-im-a-newbie/8545/7 "2009-03-30T12:29:06Z")

</div>

Oh dear - I guess we all get spoiled using decent optimizing compilers these days - time to dust off all those low level C programming tricks again then…

---

<div class="post-metadata">

**Author:** ![E.D\_Riedijk](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@E.D\_Riedijk](https://forums.developer.nvidia.com/u/E.D_Riedijk)\
**Post date:** [March 31, 2009, 12:55am UTC](https://forums.developer.nvidia.com/t/two-questions-im-a-newbie/8545/8 "2009-03-31T00:55:45Z")

</div>

> [@](#):
>
> Oh dear - I guess we all get spoiled using decent optimizing compilers these days - time to dust off all those low level C programming tricks again then…

Only if you are compute bound. I have cycles left to burn in 99% of my kernels.

---

<div class="post-metadata">

**Author:** ![Paul\_Russell](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@Paul\_Russell](https://forums.developer.nvidia.com/u/Paul_Russell)\
**Post date:** [March 31, 2009, 1:16pm UTC](https://forums.developer.nvidia.com/t/two-questions-im-a-newbie/8545/9 "2009-03-31T13:16:29Z")

</div>

> [@](#):
>
> Only if you are compute bound. I have cycles left to burn in 99% of my kernels.

Well if those cycles are going spare I could certainly use some. ;-)

---

<div class="post-metadata">

**Author:** ![E.D\_Riedijk](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@E.D\_Riedijk](https://forums.developer.nvidia.com/u/E.D_Riedijk)\
**Post date:** [March 31, 2009, 6:52pm UTC](https://forums.developer.nvidia.com/t/two-questions-im-a-newbie/8545/10 "2009-03-31T18:52:03Z")

</div>

You make me jealous :D