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 > [email protected] > http://lists.cs.uiuc.edu/mailman/listinfo/cfe-commits _______________________________________________ cfe-commits mailing list [email protected] http://lists.cs.uiuc.edu/mailman/listinfo/cfe-commits
