Integer Overflow - AMD and Intel difference

[Stephen_H_4]

Hi,

This is my first OpenCL project so I might be doing something wrong but I am having a very strange issue. 

Below is my kernel code: 

__kernel void collatz(__global int* in, __global int* out)
{
    uint id = get_global_id(0);
    unsigned long n = (unsigned long)id;
    uint count = 0;

    while (n > 1) { 
        if (n % 2 == 0) {
            n = n / 2; 
        } else { 
            if(n == 1572066143) {
                printf("BEFORE - %lu\n", n);
                n = (3 * n) + 1; 
                printf("AFTER  - %lu\n", n);
             } else {
                 n = (3 * n) + 1; 
            }

       }

       count = count + 1;
    }

    out[id] = count;
}

And here is the output:

BEFORE - 1572066143
AFTER  - 421231134

To me it looks as if the "n" variable is overflowing but I can't figure out why that is happening as it is a unsigned long. 

The really strange thing is if I update the kernel to set the value of 'n' to 1572066143 then it works correctly.

__kernel void collatz(__global int* in, __global int* out)
{
    uint id = get_global_id(0);
    unsigned long n = (unsigned long)id;
    uint count = 0;

    while (n > 1) { 
        if (n % 2 == 0) {
            n = n / 2; 
        } else { 
            if(n == 1572066143) {
		n = 1572066143;
                printf("BEFORE - %lu\n", n);
                n = (3 * n) + 1; 
                printf("AFTER  - %lu\n", n);
             } else {
                 n = (3 * n) + 1; 
            }

       }

       count = count + 1;
    }

    out[id] = count;
}

Output:

BEFORE - 1572066143
AFTER  - 4716198430

In addition if I run the exact same code on my AMD GPU then I get the correct value.

Please let us know what Processor, Operating System, Graphics Driver Version, and Tool Version you are using:

Processor: i5 4670K (HD Graphics 4600)
Operating System: Windows 10
Graphics Driver Version: 20.19.15.4300
Tool Version: Intel SDK for OpenCL 5.3.0.713

I have attached a zip file containing the host and kernel code. 

Any assistance would be great!

Thanks,
Stephen

1. If code is involved, it is great to create a small "Reproducer" sample and attach it to the message in the form of a zip file

[Robert_I_Intel]

Hi Stephen,

This looks like a bug in the compiler. Thanks for reporting! I will let the driver people know.