1 files changed, 384 insertions, 0 deletions
diff --git a/task_host_utilities.cu b/task_host_utilities.cu
new file mode 100644
index 0000000..e35292f
--- /dev/null
+++ b/task_host_utilities.cu
@@ -0,0 +1,384 @@
+// This file contains the implementation of the functions defined in
+// task_host_utilities.h--used by task_host_utilities.c to work with the GPU.
+#include <cuda_runtime.h>
+#include <errno.h>
+#include <stdint.h>
+#include <stdio.h>
+#include <string.h>
+#include <sys/mman.h>
+#include <sys/types.h>
+#include <sys/wait.h>
+#include <time.h>
+#include <unistd.h>
+// The number of GPU nanoseconds to spin for GetGPUTimerScale. Increasing this
+// will both increase the accuracy and the time the function takes to return.
+#define TIMER_SPIN_DURATION (2ull * 1000 * 1000 * 1000)
+// This macro takes a cudaError_t value. It prints an error message and returns
+// 0 if the cudaError_t isn't cudaSuccess. Otherwise, it returns nonzero.
+#define CheckCUDAError(val) (InternalCUDAErrorCheck((val), #val, __FILE__, __LINE__))
+// Prints an error message and returns 0 if the given CUDA result is an error.
+static int InternalCUDAErrorCheck(cudaError_t result, const char *fn,
+    const char *file, int line) {
+  if (result == cudaSuccess) return 1;
+  printf("CUDA error %d: %s. In %s, line %d (%s)\n", (int) result,
+    cudaGetErrorString(result), file, line, fn);
+  return 0;
+}
+static double CurrentSeconds(void) {
+  struct timespec ts;
+  if (clock_gettime(CLOCK_MONOTONIC_RAW, &ts) != 0) {
+    printf("Error getting time.\n");
+    exit(1);
+  }
+  return ((double) ts.tv_sec) + (((double) ts.tv_nsec) / 1e9);
+}
+// GlobalTimer64: Get 64-bit counter of current time on GPU
+// ***This is duplicated in benchmark_gpu_utilities.h***
+#if __CUDA_ARCH__ >= 300 // Kepler+
+// Returns the value of CUDA's global nanosecond timer.
+// Starting with sm_30, `globaltimer64` was added
+// Tuned for sm_5X architectures, but tested to still be efficient on sm_7X
+static __device__ inline uint64_t GlobalTimer64(void) {
+  uint32_t lo_bits, hi_bits, hi_bits_2;
+  uint64_t ret;
+  // Upper bits may rollover between our 1st and 2nd read
+  // (The bug seems constrained to certain old Jetson boards, so this
+  // workaround could probably be gated to only those GPUs.)
+  asm volatile("mov.u32 %0, %%globaltimer_hi;" : "=r"(hi_bits));
+  asm volatile("mov.u32 %0, %%globaltimer_lo;" : "=r"(lo_bits));
+  asm volatile("mov.u32 %0, %%globaltimer_hi;" : "=r"(hi_bits_2));
+  // If upper bits rolled over, lo_bits = 0
+  lo_bits = (hi_bits != hi_bits_2) ? 0 : lo_bits;
+  // SASS on older architectures (such as sm_52) is natively 32-bit, so the
+  // following three lines get optimized out.
+  ret = hi_bits_2;
+  ret <<= 32;
+  ret |= lo_bits;
+  return ret;
+}
+#elif __CUDA_ARCH__ < 200 // Tesla
+// On sm_13, we /only/ have `clock` (32 bits)
+static __device__ inline uint64_t GlobalTimer64(void) {
+  uint32_t lo_bits;
+  uint64_t ret;
+  asm volatile("mov.u32 %0, %%clock;" : "=r"(lo_bits));
+  ret = 0;
+  ret |= lo_bits;
+  return ret;
+}
+#else
+// Could use clock64 for sm_2x (Fermi), but that's untested
+#error Fermi-based GPUs (sm_2x) are unsupported!
+#endif
+// A simple kernel which writes the value of the globaltimer64 register to a
+// location in device memory.
+static __global__ void GetTime(uint64_t *time, volatile uint32_t *ready_barrier,
+    volatile uint32_t *start_barrier, volatile uint32_t *end_barrier) {
+  *ready_barrier = 1;
+  while (!*start_barrier)
+    continue;
+  *time = GlobalTimer64();
+  *end_barrier = 1;
+}
+// Allocates a private shared memory buffer containing the given number of
+// bytes. Can be freed by using FreeSharedBuffer. Returns NULL on error.
+// Initializes the buffer to contain 0.
+static void* AllocateSharedBuffer(size_t size) {
+  void *to_return = mmap(NULL, size, PROT_READ | PROT_WRITE, MAP_ANONYMOUS |
+    MAP_SHARED, -1, 0);
+  if (to_return == MAP_FAILED) return NULL;
+  memset(to_return, 0, size);
+  return to_return;
+}
+// Frees a shared buffer returned by AllocateSharedBuffer.
+static void FreeSharedBuffer(void *buffer, size_t size) {
+  munmap(buffer, size);
+}
+// This function should be run in a separate process in order to read the GPU's
+// nanosecond counter. Sets times to 0 on error.
+static void InternalReadGPUNanoseconds(int cuda_device, double *cpu_time,
+    uint64_t *gpu_time) {
+  uint64_t *device_time = NULL;
+  volatile uint32_t *gpu_start_barrier, *start_barrier = NULL;
+  volatile uint32_t *gpu_end_barrier, *end_barrier = NULL;
+  volatile uint32_t *gpu_ready_barrier, *ready_barrier = NULL;
+  volatile double cpu_start, cpu_end;
+  double max_error;
+  // Times need to be zero in case any of the following logic fails
+  *cpu_time = 0;
+  *gpu_time = 0;
+  if (!CheckCUDAError(cudaSetDevice(cuda_device))) return;
+  if (!CheckCUDAError(cudaMalloc(&device_time, sizeof(*device_time)))) return;
+  if (!CheckCUDAError(cudaHostAlloc(&start_barrier, sizeof(*start_barrier),
+    cudaHostAllocMapped))) goto out;
+  if (!CheckCUDAError(cudaHostAlloc(&end_barrier, sizeof(*end_barrier),
+    cudaHostAllocMapped))) goto out;
+  if (!CheckCUDAError(cudaHostAlloc(&ready_barrier, sizeof(*ready_barrier),
+    cudaHostAllocMapped))) goto out;
+  // Setup device pointers for all the barriers
+  if (!CheckCUDAError(cudaHostGetDevicePointer((uint32_t**)&gpu_start_barrier,
+    (uint32_t*)start_barrier, 0))) goto out;
+  if (!CheckCUDAError(cudaHostGetDevicePointer((uint32_t**)&gpu_end_barrier,
+    (uint32_t*)end_barrier, 0))) goto out;
+  if (!CheckCUDAError(cudaHostGetDevicePointer((uint32_t**)&gpu_ready_barrier,
+    (uint32_t*)ready_barrier, 0))) goto out;
+  /* Clock Synchronization Flow
+                            Start Barrier--+
+                                           |
+    CPU: Launch Kernel......^-->CPU TS 1-->|............^-->CPU TS 2-->Done
+                |           |              |            |
+    GPU:        +---------->|..............v-->GPU TS-->|-->Terminate
+                            |                           |
+             Ready Barrier--+              End Barrier--+
+    GPU TS ~= CPU TS 1 + (CPU TS 2 - CPU TS 1) / 2.0
+    If it takes about as long for a CPU memory write to DRAM to be visible to
+    the GPU as it takes for a GPU memory write to DRAM to be visible to the
+    CPU.
+    If not, the worst possible error is if one of the above legs is instant,
+    and the other is extremely slow, in which case the time estimate is off by
+    as much as half the CPU interval [(CPU TS 2 - CPU TS 1) / 2.0]. (Depending
+    on the platform's CPU cache configuration, an inbalance may be expected.)
+    The typical error should be about the difference in the amount of time it
+    takes to read the CPU TS counter and the GPU TS counter. On recent
+    platforms, the difference shouldn't be more than double-digit nanoseconds.
+  */
+  // Run the kernel a first time to warm up the GPU.
+  GetTime<<<1, 1>>>(device_time, gpu_ready_barrier, gpu_start_barrier,
+    gpu_end_barrier);
+  *start_barrier = 1;
+  if (!CheckCUDAError(cudaDeviceSynchronize())) goto out;
+  // Now run the actual time-checking kernel.
+  *start_barrier = 0;
+  *end_barrier = 0;
+  *ready_barrier = 0;
+  GetTime<<<1, 1>>>(device_time, gpu_ready_barrier, gpu_start_barrier,
+    gpu_end_barrier);
+  // Wait for kernel to initialize
+  while (!*ready_barrier)
+    continue;
+  // Immediately record CPU time and tell GPU kernel to record time
+  cpu_start = CurrentSeconds();
+  *start_barrier = 1;
+  // Wait for kernel to finish recording time, and immediately record CPU time again
+  while (!*end_barrier)
+    continue;
+  cpu_end = CurrentSeconds();
+  // GPU clock should have been stored about half-way between the CPU reads
+  *cpu_time = (cpu_end - cpu_start) / 2.0 + cpu_start;
+  if (!CheckCUDAError(cudaMemcpy(gpu_time, device_time, sizeof(device_time),
+    cudaMemcpyDeviceToHost))) goto out;
+  max_error = (cpu_end - cpu_start) / 2.0;
+  fprintf(stderr, "Time synchronized to a maximum error of +/- %f us.\n",
+    max_error * (1000.0 * 1000.0));
+out:
+  cudaFree(device_time);
+  cudaFree((void*)start_barrier);
+  cudaFree((void*)end_barrier);
+  cudaFree((void*)ready_barrier);
+}
+int GetHostDeviceTimeOffset(int cuda_device, double *host_seconds,
+  uint64_t *gpu_nanoseconds) {
+  uint64_t *shared_gpu_time = NULL;
+  double *shared_cpu_time = NULL;
+  int status;
+  pid_t pid = -1;
+  shared_gpu_time = (uint64_t *) AllocateSharedBuffer(
+    sizeof(*shared_gpu_time));
+  if (!shared_gpu_time) {
+    printf("Failed allocating shared buffer for IPC.\n");
+    return 0;
+  }
+  shared_cpu_time = (double *) AllocateSharedBuffer(sizeof(*shared_cpu_time));
+  if (!shared_cpu_time) {
+    printf("Failed allocating shared CPU time buffer for IPC.\n");
+    FreeSharedBuffer(shared_gpu_time, sizeof(*shared_gpu_time));
+    return 0;
+  }
+  pid = fork();
+  if (pid < 0) {
+    printf("Failed creating a child process to get GPU time: %s\n", strerror(
+      errno));
+    return 0;
+  }
+  if (pid == 0) {

diff --git a/task_host_utilities.cu b/task_host_utilities.cu new file mode 100644 index 0000000..e35292f --- /dev/null +++ b/task_host_utilities.cu
@@ -0,0 +1,384 @@
	1	// This file contains the implementation of the functions defined in
	2	// task_host_utilities.h--used by task_host_utilities.c to work with the GPU.
	3	#include <cuda_runtime.h>
	4	#include <errno.h>
	5	#include <stdint.h>
	6	#include <stdio.h>
	7	#include <string.h>
	8	#include <sys/mman.h>
	9	#include <sys/types.h>
	10	#include <sys/wait.h>
	11	#include <time.h>
	12	#include <unistd.h>
	13
	14	// The number of GPU nanoseconds to spin for GetGPUTimerScale. Increasing this
	15	// will both increase the accuracy and the time the function takes to return.
	16	#define TIMER_SPIN_DURATION (2ull * 1000 * 1000 * 1000)
	17
	18	// This macro takes a cudaError_t value. It prints an error message and returns
	19	// 0 if the cudaError_t isn't cudaSuccess. Otherwise, it returns nonzero.
	20	#define CheckCUDAError(val) (InternalCUDAErrorCheck((val), #val, __FILE__, __LINE__))
	21
	22	// Prints an error message and returns 0 if the given CUDA result is an error.
	23	static int InternalCUDAErrorCheck(cudaError_t result, const char *fn,
	24	const char *file, int line) {
	25	if (result == cudaSuccess) return 1;
	26	printf("CUDA error %d: %s. In %s, line %d (%s)\n", (int) result,
	27	cudaGetErrorString(result), file, line, fn);
	28	return 0;
	29	}
	30
	31	static double CurrentSeconds(void) {
	32	struct timespec ts;
	33	if (clock_gettime(CLOCK_MONOTONIC_RAW, &ts) != 0) {
	34	printf("Error getting time.\n");
	35	exit(1);
	36	}
	37	return ((double) ts.tv_sec) + (((double) ts.tv_nsec) / 1e9);
	38	}
	39
	40	// GlobalTimer64: Get 64-bit counter of current time on GPU
	41	// *This is duplicated in benchmark_gpu_utilities.h*
	42	#if __CUDA_ARCH__ >= 300 // Kepler+
	43	// Returns the value of CUDA's global nanosecond timer.
	44	// Starting with sm_30, `globaltimer64` was added
	45	// Tuned for sm_5X architectures, but tested to still be efficient on sm_7X
	46	static __device__ inline uint64_t GlobalTimer64(void) {
	47	uint32_t lo_bits, hi_bits, hi_bits_2;
	48	uint64_t ret;
	49	// Upper bits may rollover between our 1st and 2nd read
	50	// (The bug seems constrained to certain old Jetson boards, so this
	51	// workaround could probably be gated to only those GPUs.)
	52	asm volatile("mov.u32 %0, %%globaltimer_hi;" : "=r"(hi_bits));
	53	asm volatile("mov.u32 %0, %%globaltimer_lo;" : "=r"(lo_bits));
	54	asm volatile("mov.u32 %0, %%globaltimer_hi;" : "=r"(hi_bits_2));
	55	// If upper bits rolled over, lo_bits = 0
	56	lo_bits = (hi_bits != hi_bits_2) ? 0 : lo_bits;
	57	// SASS on older architectures (such as sm_52) is natively 32-bit, so the
	58	// following three lines get optimized out.
	59	ret = hi_bits_2;
	60	ret <<= 32;
	61	ret \|= lo_bits;
	62	return ret;
	63	}
	64	#elif __CUDA_ARCH__ < 200 // Tesla
	65	// On sm_13, we /only/ have `clock` (32 bits)
	66	static __device__ inline uint64_t GlobalTimer64(void) {
	67	uint32_t lo_bits;
	68	uint64_t ret;
	69	asm volatile("mov.u32 %0, %%clock;" : "=r"(lo_bits));
	70	ret = 0;
	71	ret \|= lo_bits;
	72	return ret;
	73	}
	74	#else
	75	// Could use clock64 for sm_2x (Fermi), but that's untested
	76	#error Fermi-based GPUs (sm_2x) are unsupported!
	77	#endif
	78
	79	// A simple kernel which writes the value of the globaltimer64 register to a
	80	// location in device memory.
	81	static __global__ void GetTime(uint64_t time, volatile uint32_t ready_barrier,
	82	volatile uint32_t start_barrier, volatile uint32_t end_barrier) {
	83	*ready_barrier = 1;
	84	while (!*start_barrier)
	85	continue;
	86	*time = GlobalTimer64();
	87	*end_barrier = 1;
	88	}
	89
	90	// Allocates a private shared memory buffer containing the given number of
	91	// bytes. Can be freed by using FreeSharedBuffer. Returns NULL on error.
	92	// Initializes the buffer to contain 0.
	93	static void* AllocateSharedBuffer(size_t size) {
	94	void *to_return = mmap(NULL, size, PROT_READ \| PROT_WRITE, MAP_ANONYMOUS \|
	95	MAP_SHARED, -1, 0);
	96	if (to_return == MAP_FAILED) return NULL;
	97	memset(to_return, 0, size);
	98	return to_return;
	99	}
	100
	101	// Frees a shared buffer returned by AllocateSharedBuffer.
	102	static void FreeSharedBuffer(void *buffer, size_t size) {
	103	munmap(buffer, size);
	104	}
	105
	106	// This function should be run in a separate process in order to read the GPU's
	107	// nanosecond counter. Sets times to 0 on error.
	108	static void InternalReadGPUNanoseconds(int cuda_device, double *cpu_time,
	109	uint64_t *gpu_time) {
	110	uint64_t *device_time = NULL;
	111	volatile uint32_t gpu_start_barrier, start_barrier = NULL;
	112	volatile uint32_t gpu_end_barrier, end_barrier = NULL;
	113	volatile uint32_t gpu_ready_barrier, ready_barrier = NULL;
	114	volatile double cpu_start, cpu_end;
	115	double max_error;
	116	// Times need to be zero in case any of the following logic fails
	117	*cpu_time = 0;
	118	*gpu_time = 0;
	119	if (!CheckCUDAError(cudaSetDevice(cuda_device))) return;
	120	if (!CheckCUDAError(cudaMalloc(&device_time, sizeof(*device_time)))) return;
	121	if (!CheckCUDAError(cudaHostAlloc(&start_barrier, sizeof(*start_barrier),
	122	cudaHostAllocMapped))) goto out;
	123	if (!CheckCUDAError(cudaHostAlloc(&end_barrier, sizeof(*end_barrier),
	124	cudaHostAllocMapped))) goto out;
	125	if (!CheckCUDAError(cudaHostAlloc(&ready_barrier, sizeof(*ready_barrier),
	126	cudaHostAllocMapped))) goto out;
	127	// Setup device pointers for all the barriers
	128	if (!CheckCUDAError(cudaHostGetDevicePointer((uint32_t**)&gpu_start_barrier,
	129	(uint32_t*)start_barrier, 0))) goto out;
	130	if (!CheckCUDAError(cudaHostGetDevicePointer((uint32_t**)&gpu_end_barrier,
	131	(uint32_t*)end_barrier, 0))) goto out;
	132	if (!CheckCUDAError(cudaHostGetDevicePointer((uint32_t**)&gpu_ready_barrier,
	133	(uint32_t*)ready_barrier, 0))) goto out;
	134	/* Clock Synchronization Flow
	135
	136	Start Barrier--+
	137	\|
	138	CPU: Launch Kernel......^-->CPU TS 1-->\|............^-->CPU TS 2-->Done
	139	\| \| \| \|
	140	GPU: +---------->\|..............v-->GPU TS-->\|-->Terminate
	141	\| \|
	142	Ready Barrier--+ End Barrier--+
	143
	144	GPU TS ~= CPU TS 1 + (CPU TS 2 - CPU TS 1) / 2.0
	145
	146	If it takes about as long for a CPU memory write to DRAM to be visible to
	147	the GPU as it takes for a GPU memory write to DRAM to be visible to the
	148	CPU.
	149
	150	If not, the worst possible error is if one of the above legs is instant,
	151	and the other is extremely slow, in which case the time estimate is off by
	152	as much as half the CPU interval [(CPU TS 2 - CPU TS 1) / 2.0]. (Depending
	153	on the platform's CPU cache configuration, an inbalance may be expected.)
	154
	155	The typical error should be about the difference in the amount of time it
	156	takes to read the CPU TS counter and the GPU TS counter. On recent
	157	platforms, the difference shouldn't be more than double-digit nanoseconds.
	158	*/
	159	// Run the kernel a first time to warm up the GPU.
	160	GetTime<<<1, 1>>>(device_time, gpu_ready_barrier, gpu_start_barrier,
	161	gpu_end_barrier);
	162	*start_barrier = 1;
	163	if (!CheckCUDAError(cudaDeviceSynchronize())) goto out;
	164	// Now run the actual time-checking kernel.
	165	*start_barrier = 0;
	166	*end_barrier = 0;
	167	*ready_barrier = 0;
	168	GetTime<<<1, 1>>>(device_time, gpu_ready_barrier, gpu_start_barrier,
	169	gpu_end_barrier);
	170	// Wait for kernel to initialize
	171	while (!*ready_barrier)
	172	continue;
	173	// Immediately record CPU time and tell GPU kernel to record time
	174	cpu_start = CurrentSeconds();
	175	*start_barrier = 1;
	176	// Wait for kernel to finish recording time, and immediately record CPU time again
	177	while (!*end_barrier)
	178	continue;
	179	cpu_end = CurrentSeconds();
	180	// GPU clock should have been stored about half-way between the CPU reads
	181	*cpu_time = (cpu_end - cpu_start) / 2.0 + cpu_start;
	182	if (!CheckCUDAError(cudaMemcpy(gpu_time, device_time, sizeof(device_time),
	183	cudaMemcpyDeviceToHost))) goto out;
	184	max_error = (cpu_end - cpu_start) / 2.0;
	185	fprintf(stderr, "Time synchronized to a maximum error of +/- %f us.\n",
	186	max_error * (1000.0 * 1000.0));
	187	out:
	188	cudaFree(device_time);
	189	cudaFree((void*)start_barrier);
	190	cudaFree((void*)end_barrier);
	191	cudaFree((void*)ready_barrier);
	192	}
	193
	194	int GetHostDeviceTimeOffset(int cuda_device, double *host_seconds,
	195	uint64_t *gpu_nanoseconds) {
	196	uint64_t *shared_gpu_time = NULL;
	197	double *shared_cpu_time = NULL;
	198	int status;
	199	pid_t pid = -1;
	200	shared_gpu_time = (uint64_t *) AllocateSharedBuffer(
	201	sizeof(*shared_gpu_time));
	202	if (!shared_gpu_time) {
	203	printf("Failed allocating shared buffer for IPC.\n");
	204	return 0;
	205	}
	206	shared_cpu_time = (double ) AllocateSharedBuffer(sizeof(shared_cpu_time));
	207	if (!shared_cpu_time) {
	208	printf("Failed allocating shared CPU time buffer for IPC.\n");
	209	FreeSharedBuffer(shared_gpu_time, sizeof(*shared_gpu_time));
	210	return 0;
	211	}
	212	pid = fork();
	213	if (pid < 0) {
	214	printf("Failed creating a child process to get GPU time: %s\n", strerror(
	215	errno));
	216	return 0;
	217	}
	218	if (pid == 0) {