jjfumero / codeblogarticles Goto Github PK

Hi Juan,

Thanks for the articles and source on getting started with Level Zero.

I noticed that you time your GPU execution from zeCommandListClose to zeCommandQueueSynchronize.

I think a more accurate timing would be to measure from zeCommandQueueExecuteCommandLists with a fence to zeFenceHostSynchronize (see example here).

I've only just started to look at Level Zero and I'm not even familiar with Vulkan, but IIUC zeCommandQueueSynchronize just guarantees that the command queue has 'shipped' and is ready for reuse. It does not guarantee that execution on the GPU has completed.

I noticed that host-synchronizing on an event without synchronizing the command queue first gives an error. So I'm a little surprised about the example linked above which does not seem to require queue synchronization for fences.

Thank you for the write up. In your article, when you initialize Level Zero structures, you do not initialize certain fields. in particular, pnext as shown here:

ze_context_desc_t contextDescription = {};
contextDescription.stype = ZE_STRUCTURE_TYPE_CONTEXT_DESC;
contextDescription.flags = 0;

Intel compilers initialize them to 0 by default, whereas clang does not resulting in jumps based on uninitialized values and your samples crashing/hangning every now and then.

Fix:

ze_context_desc_t contextDescription = {};
contextDescription.stype = ZE_STRUCTURE_TYPE_CONTEXT_DESC;
contextDescription.flags = 0;
contextDescription.pNext = nullptr;

// from mxm.cpp
VALIDATECALL(zeKernelSetArgumentValue(kernel, 0, sizeof(dstResult), &dstResult));
VALIDATECALL(zeKernelSetArgumentValue(kernel, 1, sizeof(sharedA), &sharedA));
VALIDATECALL(zeKernelSetArgumentValue(kernel, 2, sizeof(sharedB), &sharedB));
VALIDATECALL(zeKernelSetArgumentValue(kernel, 3, sizeof(int), &items));

doesn't correspond to

// from matrixMultiply.cl
__kernel void mxm(__global int* a, __global int* b, __global int *c, const int n) {