[SYCL][ESIMD] Introduce rdtsc API #12315
Conversation
| {"srnd", {"srnd", {a(0), a(1)}}}}; | ||
| {"srnd", {"srnd", {a(0), a(1)}}}, | ||
| {"timestamp",{"timestamp",{}}}, | ||
| {"sr0",{"sr0",{}}}}; |
There was a problem hiding this comment.
- src0 has more info, the upper elements (2xI32 of 4) could be useful too.
- sr0 is very machine specific (return is varied depending on GPU) it gets more info in newer GPUs
let's drop support of sr0 for a while, until we know better how it is going to be used.
sycl/test-e2e/ESIMD/rdtsc_sr0.cpp
Outdated
|
|
||
| int TestResult = 0; | ||
|
|
||
| TestResult |= test_rdtsc_sr0(); | ||
|
|
||
| if (!TestResult) { | ||
| std::cout << "Pass" << std::endl; | ||
| } |
There was a problem hiding this comment.
nit: it could be shorter without extra-spaces and braces:
| int TestResult = 0; | |
| TestResult |= test_rdtsc_sr0(); | |
| if (!TestResult) { | |
| std::cout << "Pass" << std::endl; | |
| } | |
| int TestFailed = test_rdtsc_sr0(); | |
| if (!TestFailed) | |
| std::cout << "Pass" << std::endl; |
There was a problem hiding this comment.
can you please remove extra space too?
fineg74
changed the title
sycl/test-e2e/ESIMD/rdtsc.cpp
Outdated
| VectorOutputSR0.begin()); | ||
| // Check if returned values are positive | ||
| Result |= std::any_of(VectorOutputRDTSC.begin(), VectorOutputRDTSC.end(), | ||
| [](uint64_t v) { return v <= 0; }); |
There was a problem hiding this comment.
if v is unsigned, then how it can be less than zero?
The VectorOutputRDTSC needs to be signed.
There was a problem hiding this comment.
If it will be signed there will be risk of overflow. I reworked the logic to reduce the chances of overflow
There was a problem hiding this comment.
If suppose 2Ghz, then 2 secs = 32-bit. the timestamp returns 64-bit, which is huge (2 sec * 4Bil, so, there was no risk of overflow). The updated solution is good, except it needs 1 minor fix to avoid DCE-ing the load between rdtsc calls.
sycl/test-e2e/ESIMD/rdtsc.cpp
Outdated
| uint64_t StartCounter = sycl::ext::intel::experimental::esimd::rdtsc(); | ||
| simd<uint64_t, 1> VectorResultRDTSC(VectorOutputRDTSCPtr + Idx); | ||
| uint64_t EndCounter = sycl::ext::intel::experimental::esimd::rdtsc(); | ||
| VectorResultRDTSC = EndCounter > StartCounter; |
There was a problem hiding this comment.
because the original result of the load at L45 is not used, it may be deleted.
You can update L47 this way:
| VectorResultRDTSC = EndCounter > StartCounter; | |
| VectorResultRDTSC += EndCounter > StartCounter; |
| @@ -1723,6 +1723,13 @@ __ESIMD_API std::enable_if_t<__ESIMD_DNS::is_esimd_scalar<T>::value && | |||
| return __ESIMD_NS::bfn<FuncControl>(src0, src1, src2); | |||
| } | |||
|
|
|||
| /// rdtsc - get the value of rdtsc register. | |||
There was a problem hiding this comment.
rdtsc stands for ReaD TimeStamp Counter (I think). Can you please make it more clear in the comment here?
There was a problem hiding this comment.
GenX returns 4 integers (which is probably the content of the drtsc register (if it is a register).
The function only returns uint64 (part of what GenX can return), please don't use the word register in this decription.
No description provided.