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

[Colin_Riddell]

Hi all,

Working with the OpenCL 1.2 CTS I discovered one of the tests fails with the tip revision of clang. Firstly, I wonder if anyone else working with OpenCL has came across an assertion like this:

Type.cpp:841: clang::Type::ScalarTypeKind clang::Type::getScalarTypeKind() const: Assertion `isScalarType()' failed.

As far as I can tell, non scalar values are trying to be implicitly casted where vector types should be allowed to be arithmetically converted to the element type used by the vector operand as well.

Caused by CL code like:

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

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

Secondly, I would like to suggest a patch for review. I have included the above example as a lit test in my patch.

Thanks,

vector_ternary_conv.patch (1.53 KB)

[ssahasra]

I can see the crash too. The problem is that the call to checkConditionalConvertScalarsToVectors() has already implicitly casted LHS and RHS to the same vector type as CondTy, but the subsequent code still assumes scalar types. But I am confused about this conversion. The OpenCL 1.2 spec says on page 220 that "The second and third expressions can be any type, as long their types match, or there is a conversion ... <snip>". 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. Then what is the type of the following expression?

float2 = char2 ? float2 : float2

I actually get an error for this: "error: initializing 'float2' (vector of 2 'float' values) with an expression of incompatible type 'char2' (vector of 2 'char' values)". It seems Clang decided that the ternary expression has type char2. But from my understanding of the spec, the type should be float2.

How about the following? Are they valid, and what is their return type?

char2 ? int : float
char2 ? int : unsigned
float2 ? char : unsigned

Sameer.

[Colin_Riddell]

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.

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?

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.

Colin

[ssahasra]

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 that can be applied to one of the expressions to make their types match. 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.

ternary-test2.cl (618 Bytes)

[Colin_Riddell]

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.

Agreed, and thanks for the test. Added that to the patch.

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.

I think we are on the same track. I was not talking about promoting both LHS and RHS, just one of them.

It's worth pointing out that the select generation is handled in CGExprScalar.cpp Value *ScalarExprEmitter::
VisitAbstractConditionalOperator()

Attached is another patch for review, including both tests. Now ignore the first patch I submitted. This new change passes all the tests run with check-clang.

Regards,
Colin

vector_ternary_conv.patch (2.55 KB)

[ssahasra]

Erm … turns out that my test is wrong for a different reason! The expression “char2 ? float : float” is not compatible with the select builtin because the number of bits in the element types needs to be the same. I am guessing that this bitness issue might have been the original inspiration for casting to CondTy instead of ResTy, that you fixed. But your fix is correct within its context. Thanks! That actually saved me from embarking on a spelunking expedition! There could be more to this than the existing tests. Please see the attached document where I have described my understanding of how the ternary operator is supposed to work. Would appreciate any faults that you find! I hope it is not overkill, but the effort was already useful — it helped me discover that my testcase above is invalid. Another example which we discussed in previous emails is also wrong: “float ? char : unsigned”, because the condition cannot be a floating point type. Sameer.

ternary-operator.txt (5.58 KB)

[ssahasra]

One issue with the patch. It should replace CondTy only in one place, not both. We still need to check if CondTy is a vector. This definitely needs more tests! Sameer.

[Colin_Riddell]

Hi Sameer,

Yes, this is a good exercise for me as well. It's really forcing me to question my understanding, rather than just hacking something together to make it work. I agree we need more tests, so I will work on a larger set of both positive and negative tests once we have solidified our understanding.

I reviewed your document and moved it around slightly to fit with my understanding of the spec. In summary, my changes are:

  * Select rules about bits can be simplified down to something very succinct like: 'use select when exp1 result is a vector. When using select, exp2 and exp3 must match.' (and conversion rules apply to make them match detailed in 4.2). This is because, beyond checking these types, at this stage in the compilation we have no other control.. which is what the checkConditionalConvertScalarsToVectors() function is attempting to do (though this may be incorrect).

I think these changes are more around the order and organization of the rules, rather than the actual rules themselves and think we are starting to come to agreement on them. Let me know what you think of my changes to your document..

Regards,
Colin

ternary-operator2.txt (2.8 KB)

[ssahasra]

Good to know that. It definitely affects AMD’s frontend and I am guessing that it also affects the Khronos SPIR generator. Thanks for looking into it! I have replied below by copying relevant parts from the file. Sure, but one motivation for writing this document was to identify gray areas, where Clang is following conventions in the absence of rules. See below regarding point (6) from my original writeup. Pasting some of the lines from your attached file: We don’t actually need to say all this. It’s enough to just cite Section 6.2.6. This was my intention in point (5): Aside: Does the minus (-) sign mean that you think this should be deleted? Also, I think that Section 6.2.1 is not really necessary in this situation, the way it occurs on page 220 of the spec: … or there is a conversion in section 6.2.1 Implicit Conversions that can be applied to one of the expressions to make their types match, … Which expression should be implicitly converted? As far as I can make out, there was no need to mention Section 6.2.1 at this point in the spec itself, especially because Section 6.2.6 is meant to clarify exactly this choice: “The purpose is to determine a common real type for the operands and result.”. This section cites the C99 spec when both operands are scalars: “Otherwise, if all operands are scalar, the usual arithmetic conversions apply, per section 6.3.1.8 of the C99 standard.” Yeah, if we have already said that the vector operator behaves like the select function call, then it implies that both are evaluated. But I thought it would be good to make this an explicit point because it was the first thing that tripped me up when I first encountered the ternary operator long ago! This is an important point; it covers exactly the part that caused the original bug. By “result type”, I meant the one obtained by usual arithmetic conversions on LHS and RHS. It is the same as ResTy in the code. If we had stopped at saying that the ternary operator behaves like the select function, then it would be illegal to have an expression like “int2 ? float : float”. The vector expansion required is not specified anywhere — we want to promote ResTy (scalar float) to a vector that matches the length of CondTy, but we do not want to change its element type. Going by the letter of the spec, it is difficult to say whether the situation itself is illegal or undefined. Note that a similar situation with the select builtin will result in a signature mismatch. Clang currently implements the first version, but at least you and I agree that the second version is more useful. There is no reason to say that one behaviour is correct and the other is not. Sameer.

[Colin_Riddell]

    ... or there is a conversion in section 6.2.1 Implicit Conversions that
    can be applied to one of the expressions to make their types match, ...

Which expression should be implicitly converted? As far as I can make out, there was no need to mention Section 6.2.1 at this point in the spec itself, especially because Section 6.2.6 is meant to clarify exactly this choice: "The purpose is to determine a common real type for the operands and result.". This section cites the C99 spec when both operands are scalars: "Otherwise, if all operands are scalar, the usual arithmetic conversions apply, per section 6.3.1.8 of the C99 standard."

Either exp2 or exp3 should be converted - which does mean they need to be evaluated. So as far as I can see, we are on the same understanding.

My latest patch provides the following:

* Tests(there are 15 in total, now):
** Some positive tests for combinations of:
*** scalar? vector:vector
*** scalar ? scalar:vector
** Some negative tests for:
*** scalar ? vec2:vec4
*** scalar? vec4:vec4
*** float(vector/scalar) ? scalar:vector
*** vec ? float:float (actually your original float2 = char2 ? float : float example)

* Code in SemaExpr::CheckConditionalOperands():
** Reject ternary expressions with float as the exp1 (CondTy) type. A new error was added for that, too.
** Convert RHSTy to LHSTy or LHSTy to RHSTy if they are not equal
** Implicit conversion to vector types for RHS+LHS if they are scalars using the original checkConditionalConvertScalarsToVectors() function.

vector_ternary_conv2.patch (16.2 KB)

[ssahasra]

More about this in the comments to the patch inlined below. I have also attached an updated version of my writeup. The intention is to document all the details of the behaviour (to be) implemented in Clang and not just type conversions. This update is mostly a cleanup with some enhancements to the table of examples in the end. Now every negative example points out the reason for the error. I have attached the tests that I have been working on. These correspond to the table of examples in the writeup. There might be overlap with your tests, but I decided to dump whatever I had so you can take a look too. I have not checked the negative tests with pristine Clang, but the following fail with your patch: 02, 03, 04, 05, 08, 09 and 10. But first we have to agree that the tests are valid! Comments inlined below. > diff --git a/lib/Sema/SemaExpr.cpp b/lib/Sema/SemaExpr.cpp > index 04497f3…e37e0c5 100644 > — a/lib/Sema/SemaExpr.cpp > +++ b/lib/Sema/SemaExpr.cpp > @@ -5817,13 +5817,27 @@ QualType Sema::CheckConditionalOperands(ExprResult &Cond, ExprResult &LHS, > QualType LHSTy = LHS.get()->getType(); > QualType RHSTy = RHS.get()->getType(); > > - // If the condition is a vector, and both operands are scalar, > - // attempt to implicity convert them to the vector type to act like the > - // built in select. (OpenCL v1.1 s6.3.i) > - if (getLangOpts().OpenCL && CondTy->isVectorType()) > - if (checkConditionalConvertScalarsToVectors(*this, LHS, RHS, CondTy)) > - return QualType(); > - > + if (getLangOpts().OpenCL) { > + if (CondTy->isVectorType()){ > + if (CondTy->getAs()->getElementType()->isFloatingType()) This needs to be checked for the scalar operator too. > + Diag(Cond.get()->getLocStart(), diag::err_use_of_float_ternary_cond) > + << CondTy; > + if (LHSTy != RHSTy){ > + if (LHSTy->isVectorType() && RHSTy->isScalarType()) > + RHS = ImpCastExprToType(RHS.get(), LHSTy, CK_IntegralCast); > + else if (LHSTy->isScalarType() && RHSTy->isVectorType()) > + LHS = ImpCastExprToType(LHS.get(), RHSTy, CK_IntegralCast); > + } This should be moved to the function “Sema::UsualArithmeticConversions” in the same file. That function implements C99, but needs to be enhanced for Section 6.2.6 from OpenCL 1.2 spec, which provides additional conversion rules for vector operands. Once we have that, we don’t need to worry about LHS and RHS; instead, we just use ResTy which is computed just before the current point in the code. > + else if (LHSTy->isScalarType() && RHSTy->isScalarType()){ If we make the above fix, then we only need to check if ResTy is scalar at this point. > + // Now attempt to implicity convert them to the vector type to act like the > + // built in select. (OpenCL v1.1 s6.3.i) > + if (checkConditionalConvertScalarsToVectors(*this, LHS, RHS, CondTy)) This is still casting the element type of the result to the element type of the condition, which is exactly what the original problem was! It should do the following instead: > + return QualType(); > + return CondTy; > + } > + } > + } > + Wondering if control should ever reach beyond this point for OpenCL! > // If both operands have arithmetic type, do the usual arithmetic conversions > // to find a common type: C99 6.5.15p3,5. > if (LHSTy->isArithmeticType() && RHSTy->isArithmeticType()) { > Sameer.

ntest01.cl (615 Bytes)

ntest02.cl (675 Bytes)

ntest03.cl (611 Bytes)

ntest04.cl (612 Bytes)

ntest05.cl (611 Bytes)

ntest06.cl (676 Bytes)

ntest07.cl (645 Bytes)

ntest08.cl (614 Bytes)

ntest09.cl (608 Bytes)

ntest10.cl (610 Bytes)

ptest.cl (2.12 KB)

ternary-conversions.txt (5.38 KB)

[ssahasra]

I made a mistake in reading the code. The vector/scalar conversions are already handled by CheckVectorOperands, which is called just before calling UsualArithmeticConversions. But the current code is immediately returning from there … instead it should update ResTy and continue with further checks. Sameer.

[Colin_Riddell]

Hi Sameer,

I replied inline with comments.

> + if (getLangOpts().OpenCL) {
> + if (CondTy->isVectorType()){
> + if (CondTy->getAs<VectorType>()->getElementType()->isFloatingType())

This needs to be checked for the scalar operator too.

Agreed, thanks.

I have went through the examples you provided and marked which rules I disagree/am unsure about

Some notable examples:

   exp1 exp2 exp3 result eval reason
  --------+-------+--------+--------+--------+--------
   char2 int int error - 7.1

Can you remind me where the wording about number of bits matching is.... I am not sure about this one because 220 also states that " The second and third expressions can be any type, as long their types match". Which to me, suggests that the char2 ? int : int example is allowed. Also in table 6.14, for select() in c ? b : a a and b need to have matching types, but c can be different as long as its integer

        char2 int2 int2 error - 4.1

Same logic applies for this.

   int2 char char error - 7.1

And this.

              int2 int2 float2 error - 5

Agree it should error, but because no implicit conversion is allowed between int2 and float2

        int2 int2 float error - 5

Not sure about this one, my understanding here is that an attempt will be made at converting the float to a vec2. Maybe I misunderstand this rule?

           int2 char2 char2 error - 4.1

Same as my first point

[ssahasra]

Page 220 also states that if the condition is a vector, then the operator is equivalent to calling the select builtin. The select builtin in table 6.14 says that c must have the same number of elements and bits as the other two. Yup. The “5” here refers to point 5, which in turn specifies that “usual arithmetic conversions” (Section 6.2.6) are used. The same section says that no implicit conversion is allowed. So we are both saying the same thing. Section 6.2.6 states that the scalar cannot be of higher rank than the vector when converting. Looking at it another way, if we managed to convert the float to a float2, it cannot be implicitly converted to int2 or vice versa. And same explanation applies. The difference between 4.1 and 7.1 is that in 4.1, the result type is already a vector type and the rules of the select builtin apply right away. But in 7.1, the result type is a scalar type, and this is probably undefined behaviour. We are choosing to define it as if the scalar was turned into a vector without changing its element type, which is not covered by any conversion rule. Sameer.