Archives Discussions

[Steveyoungs]

Compiler does not detect invalid OpenCL code

Hi,
The compiler is succesfully compiling the following kernel, but I don't think the code is valid because of OpenCL 1.1 rev 36, section 6.3 clause g

"The logical operators and (&&, or (||) operate on all scalar and vector built-in types except the built-in scalar and vector float types."

Here is the example kernel:

__kernel void test()
{
    float3 a;
    if (a.s0 && a.s1 && a.s2)
    {
    }
}

Can you please confirm this is a bug and report it to development?

Steve

[MicahVillmow]

Steveyoungs,
This is correct. What is actually happening is this:

__kernel void test()
{
float3 a;
if ((a.s0 == 0) && (a.s1 == 0) && (a.s2 == 0))
{
}
}

The comparison is implicit in your code and the result of that comparison is a boolean type, so it does not violate OpenCL.

[Steveyoungs]

Hi Micha,

Thanks for the reply. I assume you meant

if ((a.s0 != 0) && (a.s1 != 0) && (a.s2 != 0))

The Intel OpenCL compiler rejects my original code

if (a.s0 && a.s1 && a.s2)

and when I reported it to Intel they cited section 6.3 clause G of the specification as the reason. Hence my reporting it as a bug in the AMD compiler!

I'm no language lawyer so I can't comment on who is right or wrong other than to say that logically one of the compilers is wrong!

Do you have any way to get this clarified by the OpenCL working group so that we can get all compilers accepting the same code?

Many thanks for your help

Steve

[MicahVillmow]

Ahh yeah, you are right.

In order for clarification you need to go the Khronos OpenCL forums here:
http://www.khronos.org/message...rds/viewforum.php?f=28

[Steveyoungs]

Hi Micah,

I've now had clarification from someone on the OpenCL working group, via the Khronos forum you suggested.

I asked

"should the line
if (a.s0 && a.s1 && a.s2)
compile or not? [due to section 6.3, clause g of the OpenCL spec 1.1 rev 36]

Intel say the above code is not valid. AMD say it is valid."

They said

"Intel is correct here. Section 6.3, item g states that "The logical operators and (&&, or (||) operate on all scalar and vector built-in types except the built-in scalar and vector float types"."

Can you report this as a bug in the AMD OpenCL compiler please?

Thanks,

Steve

[MicahVillmow]

Steveyoungs,
After further discussing this internally, this will not be fixed. The reason is that OpenCL inherits from C and C defines this interaction very well and it is viewed as a cut&paste typo to include this restriction on the logical operators that wasn't fixed in time before the OpenCL 1.1 spec was ratified. We view this as a spec error and will see about correcting it in future versions of the spec.

[Steveyoungs]

Micah,

Many thanks for the update.

I'd originally coded my kernel as "if (a.s0 && a.s1 && a.s2)" and it worked just fine on AMD. As you say, you'd expect it to work. It was only later when I tried running my code on Intels OpenCL that this problem occurred.

I'm happy with your solution. Lets hope you can get this change passed quickly.

Steve.