GBE/libocl: Add __gen_ocl_get_timestamp() to get timestamp.

Gen provide tm0 register for intra-kernel profiling.
Here we provide an API __gen_ocl_get_timestamp() to return
the timestamp in TM.

The return type is defined as:
struct time_stamp {
  ulong tick;
  uint event;
};

'tick' is a 64bit time tick. 'event' stores a value which means
whether a tmEvent has occured (non-zero) or not (0). tmEvent includes
time-impacting event such as context switch or frequency change
since last time tm0 was read.

I add a sample in the kernels/compiler_time_stamp.cl. Hope it
would help you understand how to use it.

V2:
Introduce ir::ARFRegister to avoid directly use of nr/subnr in Gen IR.
Rename __gen_ocl_extract_reg to __gen_ocl_region.
Rename beignet_get_time_stamp to __gen_ocl_get_timestamp.

Signed-off-by: Ruiling Song <ruiling.song@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>

1 files changed, 28 insertions, 0 deletions

diff --git a/kernels/compiler_time_stamp.cl b/kernels/compiler_time_stamp.cl
new file mode 100644
index 00000000..f66da58d
--- /dev/null
+++ b/kernels/compiler_time_stamp.cl
@@ -0,0 +1,28 @@
+__kernel void
+compiler_time_stamp(__global int *src, __global int *dst)
+{
+  int i;
+  int final[16];
+  struct time_stamp t1, t2, t3;
+  t1 = __gen_ocl_get_timestamp();
+  for (i = 0; i < 16; ++i) {
+    int array[16], j;
+    for (j = 0; j < 16; ++j)
+      array[j] = get_global_id(0);
+    for (j = 0; j < src[0]; ++j)
+      array[j] = 1+src[j];
+    final[i] = array[i];
+    if(i == 7)
+      t2 = __gen_ocl_get_timestamp();
+  }
+  t3 = __gen_ocl_get_timestamp();
+  // currently printf does not support long type.
+  // printf("tmEvt %d %d %d  tmDiff %lu %lu\n", t3-t1, t2-t1);
+
+  // time_stamp.event maybe not zero, then the time diff is not accurate,
+  // because a time event occurs before the time stamp.
+  printf("tmEvt %d %d %d  tmDiff %u %u\n", t1.event, t2.event, t3.event,
+        (uint)(t3.tick-t1.tick), (uint)(t2.tick-t1.tick));
+
+  dst[get_global_id(0)] = final[get_global_id(0)];
+}