ircd::simt: Consolidate timestamp counter sampling into inline.

include/ircd/simt/cycles.h

@@ -0,0 +1,26 @@
+// Matrix Construct
+//
+// Copyright (C) Matrix Construct Developers, Authors & Contributors
+// Copyright (C) 2016-2022 Jason Volk <jason@zemos.net>
+//
+// Permission to use, copy, modify, and/or distribute this software for any
+// purpose with or without fee is hereby granted, provided that the above
+// copyright notice and this permission notice is present in all copies. The
+// full license for this software is available in the LICENSE file.
+
+#pragma once
+#define HAVE_IRCD_SIMT_CYCLES_H
+
+#if defined(__OPENCL_VERSION__)
+inline ulong
+__attribute__((always_inline))
+ircd_simt_cycles()
+{
+	// Compiles but doesn't link on SPIR
+	#if __has_builtin(__builtin_readcyclecounter) && !defined(__SPIR)
+		return __builtin_readcyclecounter();
+	#else
+		return 0;
+	#endif
+}
+#endif

include/ircd/simt/simt.h

@@ -13,6 +13,7 @@
 
 #include "portable.h"
 #include "assert.h"
+#include "cycles.h"
 #include "broadcast.h"
 #include "reduce_add.h"
 #include "reduce_max.h"

ircd/gpt_gpu.cl

@@ -56,10 +56,8 @@ ircd_gpt_enter(__global const void *const restrict model,
 	ln = get_local_size(0),
 	cycle = ctrl->clk.cycle;
 
-	#if defined(__clang__) && !defined(__SPIR)
 	if(li == 0)
-		ctrl->prof.entered = __builtin_readcyclecounter();
-	#endif
+		ctrl->prof.entered = ircd_simt_cycles();
 }
 
 __kernel void
@@ -1007,10 +1005,8 @@ ircd_gpt_leave(__global const void *const restrict model,
 	stepping = sampling && (ctrl->clk.samp + 1) >= batch_size,
 	epoching = stepping && (ctrl->clk.step + 1) >= steps;
 
-	#if defined(__clang__) && !defined(__SPIR)
 	if(li == 0)
-		ctrl->prof.finished = __builtin_readcyclecounter();
-	#endif
+		ctrl->prof.finished = ircd_simt_cycles();
 
 	if(li == 0)
 	{