const uint4 assignment to images

Accelerated Computing CUDA CUDA Programming and Performance
1

This is a nearly trivial kernel from an implementation of John Conway’s Game of Life.

Each pixel reads its neighbors and sets itself to either cyan or white. It works fine, as is.

The catch: I can’t set the color to anything known at compile time. It generates a runtime

error. This happens only inside a conditional. See the commented line below.

Any ideas?

-Robert

[codebox]

const sampler_t mysampler =

  CLK_NORMALIZED_COORDS_FALSE|CLK_ADDRESS_CLAMP_TO_EDGE|CLK_FI

LTER_NEAREST;

__kernel void flip(__read_only image2d_t iimage, __write_only image2d_t oimage)

{

unsigned int x = get_global_id(0);

unsigned int y = get_global_id(1);

uint4 mycolor, myvalue;

const uint4 cyan = (uint4)(0,128,128,0);

const uint4 white = (uint4)(255,255,255,0);

mycolor = read_imageui(iimage,mysampler,(int2)(x,y));

myvalue = read_imageui(iimage,mysampler,(int2)(x,y-1));

myvalue += read_imageui(iimage,mysampler,(int2)(x,y+1));

myvalue += read_imageui(iimage,mysampler,(int2)(x-1,y));

myvalue += read_imageui(iimage,mysampler,(int2)(x+1,y));

myvalue += read_imageui(iimage,mysampler,(int2)(x-1,y-1));

myvalue += read_imageui(iimage,mysampler,(int2)(x+1,y-1));

myvalue += read_imageui(iimage,mysampler,(int2)(x-1,y+1));

myvalue += read_imageui(iimage,mysampler,(int2)(x+1,y+1));

// A change of this assignment to anything known at compile time,

// e.g., mycolor = cyan; or mycolor=(uint4)(cyan.xyz,0); will generate an error.

if(myvalue.x < 2255 || myvalue.x > 3255) mycolor=(uint4)(cyan.xyz,myvalue.w);

if(myvalue.x == 3*255) mycolor = (uint4)(white.xyz,myvalue.w);

write_imageui(oimage,(int2)(x,y),mycolor);

}

[/codebox]

2

Graduate student, Meng Zhu, has proposed an entertaining alternative:
use the oxymoronic declaration

volatile const uint4 cyan = (uint4)(0,128,128,0);

and then the assignment, mycolor = cyan, works fine inside the conditional.
Of course, this is in effect the same solution, but why should it be necessary?

-Robert