Clarify the definition of prerequisites to commands

[kpet]

• Reword the first source of prerequisites so the wording is symmetrical with respect to the others (i.e. the first/second/third ...).
• Broaden the first source of prerequisites to cover all implicit dependencies and provide an exhaustive list of how they arise: either because of barriers or because of ordering in in-order command-queues.

Change-Id: Ic464066261fe13756347bafb4878cd6ffb5a8427

[bashbaug]

This is one of the more confusing parts of the OpenCL spec, and this isn't common, but strictly speaking I don't think this is true: if a barrier is in an out-of-order queue, and if the barrier has an explicit event wait list, then the barrier will only wait for the commands described by the event wait list commands to complete, and subsequent commands will only wait for the barrier to complete. This means that other commands enqueued before the barrier, but that the barrier did not wait on, may still be executing.

The easiest way to describe this may be to describe dependency on the barrier itself, not on the commands before the barrier, something like:

• Commands enqueued after a command-queue barrier have an implicit dependency on the command-queue barrier as a prerequesite.

Or, slightly more verbose:

• Commands enqueued after a command-queue barrier have an implicit dependency on the completion of the preceding command-queue barrier as a prerequesite.

[kpet]

I agree with your description of what the behaviour should be. This was meant to cover barriers with no explicit prerequisites, good catch. I've updated the description.

I don't think your proposed "dependency on the barrier" wording (paraphrasing) works well. We'd still need to describe what a barrier depends on and that is only done in the description for clEnqueueBarrierWithWaitList and not in terms of prerequisites.

[bashbaug]

I'm fine with the updated text and this is a definite improvement over what is currently in the spec, but FWIW I do quite like the "dependency on the barrier" wording. Consider a case like the following (pseudocode):

clEnqueueNDRangeKernel(A);
clEnqueueNDRangeKernel(B);
clEnqueueNDRangeKernel(C);
clEnqueueBarrier(wait on {A, B});
clEnqueueNDRangeKernel(D);

Do we agree:

1. Kernels A, B, and C have no implicit or explicit prerequisites, so they may execute anytime and in any order, and may overlap execution.
2. The barrier will not complete until kernels A and B are complete because of the explicit dependencies specified by the event wait list.
3. Kernel D has the barrier as an implicit prerequisite, hence a transitive dependency on kernels A and B, even though it has no explicit prerequisites. There is no dependency between kernels C and D though, so C and D may overlap execution.

If so, I think (2) is covered by "the second source of prerequisites", but I'm not sure how (3) would be covered by the text in this PR.

[kpet]

1. Yes.
2. Yes.
3. Yes.

I've returned to this with a fresher mind and definitely agree your suggestion of describing implicit prerequisites in terms on dependencies on barriers is better. I've tweaked the PR to reflect that which I believe covers your case 3.
PTAL.