#include <iostream>
#include <hip/hip_runtime.h>
// Simple time-consuming kernel without arguments
__global__ void slowKernel() {
float val = 0.0f;
for (int i = 0; i < 10000; i++) {
for (int j = 0; j < 10000; j++) {
val += sqrtf(val + i + j);
}
}
}
int main() {
float milliseconds = 0;
hipEvent_t start, stop;
hipEventCreate(&start);
hipEventCreate(&stop);
hipEventRecord(start, 0);
// Launching the kernel with arbitrary grid and block sizes
hipLaunchKernelGGL(slowKernel, dim3(512), dim3(256), 0, 0);
hipEventRecord(stop, 0);
assert(hipEventElapsedTime(&milliseconds, stop, stop) == hipErrorNotReady);
assert(hipEventElapsedTime(&milliseconds, start, start) == hipSuccess);
assert(hipEventElapsedTime(&milliseconds, stop, stop) == hipErrorNotReady);
hipError_t err = hipEventElapsedTime(&milliseconds, start, stop);
// Check if elapsed time returns hipErrorNotReady
if (err == hipErrorNotReady) {
std::cout << "Kernel still in progress..." << std::endl;
} else {
std::cout << "Kernel execution time: " << milliseconds << " ms" << std::endl;
}
return 0;
}
This kernel runs for about 3 seconds, and as you can see in the example, I am not doing any explicit synchronization.
Issue: event query on the start event seems to force some sort of synchronization such that the repeated stop event query returns ZE_RESULT_SUCCESS even though the kernel should still be running. If I query the stop event multiple times, it will return ZE_NOT_READY but as soon as I query the start event, the next query to stop event will result in ZE_RESULT_SUCCESS while not ready is expected.
Here's the Level Zero log - please let me know if this is some sort of expected behavior and if not, I can make a pure level zero reproducer.
I have the following simple test case:
This kernel runs for about 3 seconds, and as you can see in the example, I am not doing any explicit synchronization.
Issue: event query on the
startevent seems to force some sort of synchronization such that the repeatedstopevent query returnsZE_RESULT_SUCCESSeven though the kernel should still be running. If I query the stop event multiple times, it will returnZE_NOT_READYbut as soon as I query thestartevent, the next query tostopevent will result inZE_RESULT_SUCCESSwhile not ready is expected.Here's the Level Zero log - please let me know if this is some sort of expected behavior and if not, I can make a pure level zero reproducer.