Strange behvior on a simple test case Result of a very simple arithmetic process is wrong

Accelerated Computing CUDA CUDA Programming and Performance
1

There is an error inside one of my openCL development. After having spent some time, i’ve found the simplest CL code that fails:

kernel void Algo( __global   int* outp )

{

	int PDif;

	int PShift;

	int PAbs;

	__local int P0;

	if(get_local_id(0)==0) {

		P0 = 2;

	}

	barrier(CLK_LOCAL_MEM_FENCE);

	

	if(get_local_id(0)==0) {

		PShift  = P0 << 1;       // Should be 4

		PDif    = 1 - P0;        // Should be -1 

		PAbs    = abs(PDif);     // Should be 1

		outp[0] = PShift + PAbs; // Should be 5

	}

}

The result should always be 5 since the output buffer value only depends on a constant value (P0) stored in a local variable.

Unfortunately, the result is … 3 ! As if the PAbs result was negative !

There are various ways to get a correct behavior such as using the -cl-optim-disable option, adding a local barrier after the computation of the PShift variable or declaring the P0 variable into the private memory.

The platform is:

device 0x25f1490 has              Quadro 2000 CL_DEVICE_NAME

device 0x25f1490 has       NVIDIA Corporation CL_DEVICE_VENDOR

device 0x25f1490 has                   295.49 CL_DRIVER_VERSION

device 0x25f1490 has          OpenCL 1.1 CUDA CL_DEVICE_VERSION

device 0x25f1490 has             FULL_PROFILE CL_DEVICE_PROFILE

device 0x25f1490 has                        4 CL_DEVICE_MAX_COMPUTE_UNITS

device 0x25f1490 has                     1024 CL_DEVICE_MAX_WORK_GROUP_SIZE

device 0x25f1490 has                    65536 CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE

device 0x25f1490 has               1072889856 CL_DEVICE_GLOBAL_MEM_SIZE

device 0x25f1490 has                    49152 CL_DEVICE_LOCAL_MEM_SIZE

device 0x25f1490 has                        2 CL_DEVICE_COMPUTE_CAPABILITY_MAJOR_NV

device 0x25f1490 has                        1 CL_DEVICE_COMPUTE_CAPABILITY_MINOR_NV

device 0x25f1490 has                    32768 CL_DEVICE_REGISTERS_PER_BLOCK_NV

[For an unknown reason, i cannot attach the host C code. If needed, i will post it]

So what is wrong with this code ?

Thanks for any help.

2

make sure you are managing your buffers right.

3

Thanks for your reply.

I think there is no problem with the way the buffers are managed.
There are several algorithms (each of several hundred CL lines) that are working on images up to 8 MPixl.
There is only one that fails. From this latter, i’ve extracted the piece of code generating the error to make this test case.
I can send you the Host C Code if you want to check it.

4

I’ve always found it instructive to output intermediate results when things go bad. Try something like this (and you may need to move the outp = statements):

<snip>

	if(get_local_id(0)==0) {

		PShift  = P0 << 1;       // Should be 4

		PDif    = 1 - P0;        // Should be -1 

		PAbs    = abs(PDif);     // Should be 1

		outp[0] = PShift + PAbs; // Should be 5

                outp[1] = PShift;

                outp[2] = PAbs;

                outp[3] = P0;

                outp[4] = abs ( -1 );

                // etc...

	}
5

Yes, you are right. It is always instructive and this is the reason why i did it.
But, unfortunetaly, tracing intermediate buffers leads to a correct behiavior … Final and intermediate values are right in case of tracing …

How can i determine if it is a compiler or an hardware error without changing the graphic board ?