# break-statement in kernels

**URL:** <https://forums.developer.nvidia.com/t/break-statement-in-kernels/1650>\
**Category:** CUDA Programming and Performance\
**Created:** [October 17, 2007, 12:19pm UTC](https://forums.developer.nvidia.com/t/break-statement-in-kernels/1650 "2007-10-17T12:19:51Z")\
**Posts on this page:** 3\
**Page:** 1

<div class="post-metadata">

**Author:** ![weigo](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@weigo](https://forums.developer.nvidia.com/u/weigo)\
**Post date:** [October 17, 2007, 12:19pm UTC](https://forums.developer.nvidia.com/t/break-statement-in-kernels/1650/1 "2007-10-17T12:19:51Z")

</div>

Here’s a kernel that behaves very strange.

If i execute the kernel pasted below, my resulting values in **out** are (0,0), (2,0), (4,0), … (0,2) … everything fine.

If i uncomment the line _pos = nextpos;_, the y component of the resulting values in **out** are set to zero.

If i replace the _break;_ statement by _iter = 1;_ everything is fine again.

Have i missed something essential here?

```auto
void __global__ kernelConverge(short2* out, short width)

{

  short2 pos;

  pos.x = (threadIdx.x % (width/2)) * 2;

  pos.y = (threadIdx.x / (width/2)) * 2;

  short2 nextpos = pos;

 unsigned char iter = 1;

  while(iter != 0)

  {

    if ((nextpos.y == pos.y) && (nextpos.x == pos.

    {

      break;

    }

   //pos = nextpos;

    iter--;

  }

  out[threadIdx.x] = pos;

}

```

---

<div class="post-metadata">

**Author:** ![wildcat4096](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@wildcat4096](https://forums.developer.nvidia.com/u/wildcat4096)\
**Post date:** [October 17, 2007, 6:34pm UTC](https://forums.developer.nvidia.com/t/break-statement-in-kernels/1650/2 "2007-10-17T18:34:14Z")

</div>

Would you be willing to post the support code that you are using to drive the kernel? In other words, a self contained compilable example?

---

<div class="post-metadata">

**Author:** ![weigo](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@weigo](https://forums.developer.nvidia.com/u/weigo)\
**Post date:** [October 23, 2007, 11:39am UTC](https://forums.developer.nvidia.com/t/break-statement-in-kernels/1650/3 "2007-10-23T11:39:10Z")

</div>

Sorry for the delay! Here’s is the full example.

The kernel should always jump out of the loop with the break statement immediately because nextpos was set to pos. The result is ok for the code pasted below. But if you uncomment line **pos = nextpos;** , the y components of the resulting short-ints are set to zero.

```auto
void __global__ kernelConverge(short2* out, short width)

{

 short2 pos;

 pos.x = (threadIdx.x % (width/2)) * 2;

 pos.y = (threadIdx.x / (width/2)) * 2;

 short2 nextpos = pos;

unsigned char iter = 1;

 while(iter != 0)

 {

   if ((nextpos.y == pos.y) && (nextpos.x == pos.x))

   {

     break;

   }

  //pos = nextpos;

   iter--;

 }

 out[threadIdx.x] = pos;

}

int main(int argc, char* argv[])

{

  const int SeekerCount = 144;

  const int Width = SeekerCount / 16;  //9

  const int ByteSize = sizeof(short2) * SeekerCount;

  short2* data = (short2*)malloc(ByteSize);

 short2* dataDev;

  if (cudaMalloc((void**)&dataDev, ByteSize) != cudaSuccess)

    exit(-1);

 dim3 blocks(1, 1, 1);

  dim3 threads(SeekerCount, 1, 1);

 kernelConverge<<<blocks, threads>>>(dataDev, Width);

 if (cudaMemcpy(data, dataDev, ByteSize, cudaMemcpyDeviceToHost) != cudaSuccess)

    exit(-1);

 cudaFree(dataDev);

 free(data);

}

```