[PATCH] D27334: [OpenCL] Ambiguous function call.

Tue Dec 6 11:06:32 PST 2016

I think we can’t avoid deviating from C/C++ implementation completely. But we should certainly try to avoid it in unnecessary situations as much as we can.

In this case it seems like the spec enforces the rules for resolving the ambiguity that seems logical to me.

Is this possibly a better way to deviate OpenCL implementation than giving a hypothetical warning and an error in the case the ambiguity wasn’t resolved?

Cheers,
Anastasia

From: Bader, Alexey [mailto:alexey.bader at intel.com]
Sent: 06 December 2016 12:17
To: Anastasia Stulova; Richard Smith; Bruno Cardoso Lopes via Phabricator; reviews+D27334+public+f2c5a66032c4c02e at reviews.llvm.org
Cc: egor.churaev at gmail.com; yaxun.liu at amd.com; cfe-commits at lists.llvm.org; nd
Subject: RE: [PATCH] D27334: [OpenCL] Ambiguous function call.

Actually OpenCL specification has the following text in “Built-in Functions” chapter (OpenCL C 2.0 rev. 33):

“ User defined OpenCL C functions, behave per C standard rules for functions (C99, TC2, Section
6.9.1). On entry to the function, the size of each variably modified parameter is evaluated and
the value of each argument expression is converted to the type of the corresponding parameter as
per usual arithmetic conversion rules described in section 6.2.6. Built-in functions described in
this section behave similarly, except that in order to avoid ambiguity between multiple forms of
the same built-in function, implicit scalar widening shall not occur. Note that some built-in
functions described in this section do have forms that operate on mixed scalar and vector types,
however.”

If I understand it correctly this text can help to resolve the overloading, since “usual arithmetic conversion rules” from section 6.2.6 rank floating point and integer data types including vector flavors. On the other hand I don’t think it’s good idea to adopt this rule, since it doesn’t follow C++ overloading rules and definitely will deviate OpenCL behavior from C or C++.

Thanks,
Alexey

From: Anastasia Stulova [mailto:Anastasia.Stulova at arm.com]
Sent: Monday, December 5, 2016 9:21 PM
To: Richard Smith <richard at metafoo.co.uk<mailto:richard at metafoo.co.uk>>; Bruno Cardoso Lopes via Phabricator <reviews at reviews.llvm.org<mailto:reviews at reviews.llvm.org>>; reviews+D27334+public+f2c5a66032c4c02e at reviews.llvm.org<mailto:reviews+D27334+public+f2c5a66032c4c02e at reviews.llvm.org>
Cc: Bader, Alexey <alexey.bader at intel.com<mailto:alexey.bader at intel.com>>; egor.churaev at gmail.com<mailto:egor.churaev at gmail.com>; yaxun.liu at amd.com<mailto:yaxun.liu at amd.com>; cfe-commits at lists.llvm.org<mailto:cfe-commits at lists.llvm.org>; nd <nd at arm.com<mailto:nd at arm.com>>
Subject: RE: [PATCH] D27334: [OpenCL] Ambiguous function call.

> Perhaps that is the problem (that there are two modes that do different things)? Could we make the double overload be present but unselectable to diagnose this problem in that mode too?

If we could resolve the overload candidate to prefer ‘int -> float’ than ‘int->double’, it would work best.

From: metafoo at gmail.com<mailto:metafoo at gmail.com> [mailto:metafoo at gmail.com] On Behalf Of Richard Smith
Sent: 05 December 2016 17:53
To: Bruno Cardoso Lopes via Phabricator; reviews+D27334+public+f2c5a66032c4c02e at reviews.llvm.org<mailto:reviews+D27334+public+f2c5a66032c4c02e at reviews.llvm.org>
Cc: alexey.bader at intel.com<mailto:alexey.bader at intel.com>; egor.churaev at gmail.com<mailto:egor.churaev at gmail.com>; yaxun.liu at amd.com<mailto:yaxun.liu at amd.com>; Anastasia Stulova; cfe-commits at lists.llvm.org<mailto:cfe-commits at lists.llvm.org>
Subject: Re: [PATCH] D27334: [OpenCL] Ambiguous function call.

On 5 Dec 2016 9:42 am, "Anastasia Stulova via Phabricator via cfe-commits" <cfe-commits at lists.llvm.org<mailto:cfe-commits at lists.llvm.org>> wrote:
Anastasia added a comment.

In https://reviews.llvm.org/D27334#612858, @bader wrote:

> In https://reviews.llvm.org/D27334#611703, @Anastasia wrote:
>
> > This change seems to modify normal C behavior again. Is there any strong motivation for doing this and if yes could it be done generically with C?
>
>
> Motivation:
>
>   // Non-portable OpenCL 1.2 code
>   __kernel void foo(global float* out) {
>     out[get_global_id(0)] = sin(get_global_id(0));
>   }
>
>
> This program compiles fine on OpenCL platform w/o doubles support

Perhaps that is the problem (that there are two modes that do different things)? Could we make the double overload be present but unselectable to diagnose this problem in that mode too?

and fails otherwise.
>  If OpenCL driver supports doubles it provides at least two versions of 'sin' built-in math function and compiler will not be able to choose the right one for 'size_t' argument.

Do you have a real example? If someone passes an integer to 'sin', I think there's a better warning we can give them :)

>  The goal of this warning is to let OpenCL developer know about potential issues with code portability.
I would argue this improves the portability much as it can also be misleading in some situations (because it refers to a potentially hypothetical problem). For example there can be builtin functions that only have a float parameter (without a double version of it). This is for example the case with read_image functions that take a float coordinate value between 0 and 1. Unfortunately this warning won't be triggered on read_image functions because there is an overload candidate with an int type of the same parameter too. But we can't exclude this situations to appear in the future or from some vendor extensions or even custom OpenCL code.

https://reviews.llvm.org/D27334

_______________________________________________
cfe-commits mailing list
cfe-commits at lists.llvm.org<mailto:cfe-commits at lists.llvm.org>
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits

--------------------------------------------------------------------
Joint Stock Company Intel A/O
Registered legal address: Krylatsky Hills Business Park,
17 Krylatskaya Str., Bldg 4, Moscow 121614,
Russian Federation

This e-mail and any attachments may contain confidential material for
the sole use of the intended recipient(s). Any review or distribution
by others is strictly prohibited. If you are not the intended
recipient, please contact the sender and delete all copies.
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20161206/d27716c8/attachment-0001.html>