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:
__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
