[cfe-dev] [PATCH] [OpenCL] Conversions for ternary operations between scalar and vector

Sahasrabuddhe, Sameer sameer.sahasrabuddhe at amd.com
Fri Dec 5 22:55:40 PST 2014


On 12/5/2014 11:20 PM, Colin Riddell wrote:
>
>> float2 = char2 ? float2 : float2 
> Should result in an error as far as my understanding of the spec' 
> goes. This is because neither explicit nor implicit conversions 
> between vector types are allowed in OpenCL 1.2 (page 209 - Section 
> 6.2.2  and page 209 -  Section 6.2.1 respectively). Oddly enough, 
> though. For me I don't see the error..
>
>> seems Clang decided that the ternary expression has type char2. But 
>> from my understanding of the spec, the type should be float2. 
>
> Yes, if anything it should be float2 - irrespective of the rules 
> against converting vec to vec.

These two seem to be contradicting each other. If the type of the 
ternary operator in the above example is float2, then there should be no 
error in assigning it to a float2!

Also my mistake about the error not showing up. The actual example that 
I tried was "float2 = char2 ? float : float". Attached the modified copy 
that I had created from your original lit test. This requires your patch 
to avoid the assertion. I believe this assignment should not cause any 
error at all.

>> It further says that the type of the entire expression is derived 
>> from the second and third expression, and nowhere does it mention the 
>> type of the first expression.
> Yes, the result is derived from either the second or third expression, 
> however this needs to be implicitly converted to the result type. The 
> type of the first expression is just used to work out whether the 
> result should be handled as vector or scalar..
>
>> char2 ? int : float
>> char2 ? int : unsigned
>> float2 ? char : unsigned 
> Are valid, as far as I can see, and their results would depend on the 
> type of result, not the types of the values on either side of the colon.
> Expanding your example by providing an assignment would mean that if 
> assigned to some vector type, the LHS or RHS values would be expanded 
> to fit the size of that vector and the types would be implicitly 
> converted to the base type of the result vector.
> For example..
> float3 = char2 ? int: float    would result in a float3 for LHS by 
> expansion of the  scalar values to vec3 and implicit conversion of the 
> int to float.  For RHS,  expansion of the scalar to vec3 and no 
> implicit conversion needed because it's already the same type.
>
> Do you agree with this?

No, that's not what I see in the spec on page 220: "as long their types 
match, or there is a conversion <snip> that can be applied to one of the 
expressions to make their types match. <snip> This resulting matching 
type is the type of the entire expression." The spec does not actually 
say what is this "resulting matching type", but I would say it is the 
one with greater rank (Section 6.2.6) between LHS or RHS (same as ResTy 
in the code).

In the same paragraph, the spec also says that if the condition type is 
a vector, the behaviour is similar to the "select" builtin. This means 
that the vector sizes of LHS and RHS after conversion need to match the 
condition size.

So for each of the expressions above, we first promote LHS or RHS (but 
never both) to the "resulting matching type", and then expand them to 
vectors that match the condition type. When we try to assign this to a 
float3, the expected error is that the vector sizes do not match. There 
is no other way to do it, because we can't expand the vec2 condition to 
a vec3.

> From thinking about this, I now see the two lines I proposed are not 
> enough in order to implement these rules correctly and fix the current 
> bug.
>
> I see what checkConditionalConvertScalarsToVectors() is trying to do 
> by converting LHS and RHS to vectors, too.   However, subsequent code 
> checks LHSTy and RHSTy, which are QualType values based on LHS and RHS 
> before their types were converted and someone has forgotten to update 
> them.
> I don't think it's enough to convert these values to vectors and let 
> the code that follows incorrectly handle the remainder of the 
> conversion and ternary logic assuming C99.
>
> I think the bug may be caused by trying to convert the LHS and RHS to 
> CondTy, rather than ResTy.

Right.

Sameer.
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/cfe-dev/attachments/20141206/d3261945/attachment.html>
-------------- next part --------------
typedef char char2 __attribute__((ext_vector_type(2)));
typedef char char3 __attribute__((ext_vector_type(3)));
typedef char char4 __attribute__((ext_vector_type(4)));
typedef char char8 __attribute__((ext_vector_type(8)));
typedef char char16 __attribute__((ext_vector_type(16)));

typedef float float2 __attribute__((ext_vector_type(2)));

__kernel void test(__global float *srcA, __global float *srcB, __global char2 *srcC, __global float2 *dst)
{
  int  tid = 1;

  float valA = srcA[ tid ];
  float valB = srcB[ tid ];
  char2 valC = srcC[ tid ];
  float2 destVal = valC ? valA : valB;
  dst[ tid ] = destVal;
}



More information about the cfe-dev mailing list