[PATCH] OpenCL 1.2 allow explicit zero-to-event_t cast

[Tom Stellard]

On Mon, Sep 29, 2014 at 10:54:00AM +0100, Fraser Cormack wrote:
> Hi,
> 
> This patch allows an explicit cast from the literal 0 to an event_t.
> The fact that you can use 0 as an event_t is in Section 6.12.10 of
> the OpenCL 1.2 specification. The use of an explicit cast is used in
> the async_copy* kernels in the 'basic' subtest of the conformance
> test suite:
> 

Hi,

The same patch has been done at least two times before.
See 
http://lists.cs.uiuc.edu/pipermail/cfe-commits/Week-of-Mon-20140217/099653.html

-Tom

> __kernel void test_fn( const __global short8 *src, __global short8 *dst,
>                        __local short8 *localBuffer, int copiesPerWorkgroup, int copiesPerWorkItem )
> {
>     int i;
>     for(i=0; i<copiesPerWorkItem; i++)
>         localBuffer[ get_local_id( 0 )*copiesPerWorkItem+i ] = (short8)(short)0;
> 
>     barrier( CLK_LOCAL_MEM_FENCE );
> 
>     for(i=0; i<copiesPerWorkItem; i++)
>         localBuffer[ get_local_id( 0 )*copiesPerWorkItem+i ] = src[ get_global_id( 0 )*copiesPerWorkItem+i ];
> 
>     barrier( CLK_LOCAL_MEM_FENCE );
> 
>     event_t event;
>     event = async_work_group_copy((__global short8*)(dst+copiesPerWorkgroup*get_group_id(0)),
>                                   (__local const short8*)localBuffer,
>                                   (size_t)copiesPerWorkgroup,
> ** HERE ** >>>>                   (event_t)0 );
>     wait_group_events( 1, &event );
> }
> 
> 
> I've also added a test for this behaviour.
> 
> Cheers,
> Fraser
> 
> -- 
> Fraser Cormack
> Compiler Developer
> Codeplay Software Ltd
> 45 York Place, Edinburgh, EH1 3HP
> Tel: 0131 466 0503
> Fax: 0131 557 6600
> Website: http://www.codeplay.com
> Twitter: https://twitter.com/codeplaysoft
> 
> This email and any attachments may contain confidential and /or privileged information and  is for use  by the addressee only. If you are not the intended recipient, please notify Codeplay Software Ltd immediately and delete the message from your computer. You may not copy or forward it,or use or disclose its contents to any other person. Any views or other information in this message which do not relate to our business are not authorized by Codeplay software Ltd, nor does this message form part of any contract unless so stated.
> As internet communications are capable of data corruption Codeplay Software Ltd does not accept any responsibility for any changes made to this message after it was sent. Please note that Codeplay Software Ltd does not accept any liability or responsibility for viruses and it is your responsibility to scan any attachments.
> Company registered in England and Wales, number: 04567874
> Registered office: 81 Linkfield Street, Redhill RH1 6BY
> 

> Index: lib/Sema/SemaCast.cpp
> ===================================================================
> --- lib/Sema/SemaCast.cpp	(revision 218598)
> +++ lib/Sema/SemaCast.cpp	(working copy)
> @@ -2217,6 +2217,18 @@
>      return;
>    }
>  
> +  //
> +  // OpenCL 1.2 spec, s6.12.10
> +  //
> +  // The event argument can also be used to associate the async_work_group_copy
> +  // with a previous async copy allowing an event to be shared by multiple
> +  // async copies; otherwise event should be zero.
> +  if (Self.getLangOpts().OpenCL && DestType->isEventT() &&
> +      SrcExpr.get()->EvaluateKnownConstInt(Self.getASTContext()) == 0) {
> +    Kind = CK_ZeroToOCLEvent;
> +    return;
> +  }
> +
>    if (!DestType->isScalarType() && !DestType->isVectorType()) {
>      const RecordType *DestRecordTy = DestType->getAs<RecordType>();
>  
> Index: test/CodeGenOpenCL/event_t.cl
> ===================================================================
> --- test/CodeGenOpenCL/event_t.cl	(revision 218598)
> +++ test/CodeGenOpenCL/event_t.cl	(working copy)
> @@ -9,4 +9,6 @@
>  // CHECK: call void @foo(%opencl.event_t* %
>    foo(0);
>  // CHECK: call void @foo(%opencl.event_t* null)
> +  foo((event_t)0);
> +// CHECK: call void @foo(%opencl.event_t* null)
>  }

> _______________________________________________
> cfe-commits mailing list
> cfe-commits at cs.uiuc.edu
> http://lists.cs.uiuc.edu/mailman/listinfo/cfe-commits