From e415e282be53a3d8264b67d22b4b15b518f71099 Mon Sep 17 00:00:00 2001 From: Joshua Bakita Date: Thu, 6 Jul 2023 11:45:59 -0400 Subject: Add preemption_logger benchmark, make target, and fix d82c60c Benchmark supports: => Spinning on-GPU until some number of preemptions have occured => Monitoring GPU scheduling by logging preempt and resume times => Logging the above as raw GPU, or converted CPU, times => Measuring launch overheads Also imports task_host_utilities.cu from cuda_scheduling_examiner as of commit 9fc0ca898ef64b82cbc3734c60405723650bc66f. This is needed for timing logic used in both preemption_logger and mon_cross_ctx_copies, thus the latter would not compile prior to this commit. Also adds a new testbench helper, and a .gitignore for binaries. --- .gitignore | 4 + Makefile | 7 +- preemption_logger.cu | 158 ++++++++++++++++++++ task_host_utilities.cu | 384 +++++++++++++++++++++++++++++++++++++++++++++++++ testbench.h | 3 + 5 files changed, 554 insertions(+), 2 deletions(-) create mode 100644 .gitignore create mode 100644 preemption_logger.cu create mode 100644 task_host_utilities.cu diff --git a/.gitignore b/.gitignore new file mode 100644 index 0000000..e589da8 --- /dev/null +++ b/.gitignore @@ -0,0 +1,4 @@ +preemption_logger +constant_cycles_kernel +copy_experiments/copy_contender +copy_experiments/mon_cross_ctx_copies diff --git a/Makefile b/Makefile index 16e17dd..9f6da4d 100644 --- a/Makefile +++ b/Makefile @@ -1,13 +1,16 @@ .PHONY: all copy_experiments clean -all: constant_cycles_kernel copy_experiments +all: constant_cycles_kernel preemption_logger copy_experiments constant_cycles_kernel: constant_cycles_kernel.cu testbench.h nvcc constant_cycles_kernel.cu -o constant_cycles_kernel +preemption_logger: task_host_utilities.cu preemption_logger.cu testbench.h + nvcc preemption_logger.cu -o preemption_logger -g + copy_experiments: $(MAKE) -C $@ clean: - rm -f constant_cycles_kernel + rm -f constant_cycles_kernel preemption_logger $(MAKE) -C copy_experiments clean diff --git a/preemption_logger.cu b/preemption_logger.cu new file mode 100644 index 0000000..667c459 --- /dev/null +++ b/preemption_logger.cu @@ -0,0 +1,158 @@ +/* Copyright 2023 Joshua Bakita + * Tool that tracks preemptions as discontinuities in GPU time and logs them. + * + * The logged intervals are output to standard out as a CSV, with each row + * containing a single interval. The first column contains the interval start + * time in nanoseconds (ns), and the second column contains the interval end + * time in ns. + * + * All progress and error messages and printed to standard error. + * + * Assumptions: + * - The GPU clock continues to tick, regardless of if this is the current + * application running. + */ +#include +#include +#include +#include +#include + +#include "task_host_utilities.cu" +#include "testbench.h" + +struct interval { + uint64_t start; + uint64_t end; +}; + +// Minimum time discontinuity which indicates a gap between intervals +// One clock tick is about 1ns +#define MIN_PREEMPT_TICKS 2*1000 // ~2us + +// Watch for discontinuities in the GPU clock, indicating intervals during +// which we we were preempted and another context ran. +// @param ivls Base address of array to log interval start and end times +// @param num_ivls How many intervals should be logged before exiting? +__global__ void loop_on_gpu(struct interval *ivls, int num_ivls) { + struct interval *curr_ivl = ivls; + uint64_t last_time = GlobalTimer64(); + if (num_ivls <= 0) + return; + curr_ivl->start = last_time; + // While no preemptions are occuring, this repeatedly obtains the clock, + // storing the value in `last_time`. When we're switched away from and back + // to, there will appear to be a large gap between `now` and `last_time`. + // We record `last_time` as the end of the last interval, and `now` as the + // start of the next interval, then resume repeatedly obtaining the clock. + + // Continue to log intervals until the pointer `curr_ivl` points to a + // location off the end of the `ivls` array. + while (curr_ivl < ivls + num_ivls) { + uint64_t now = GlobalTimer64(); + if (now > last_time + MIN_PREEMPT_TICKS) { + curr_ivl->end = last_time; + curr_ivl++; + curr_ivl->start = now; + } + last_time = now; + } +} + +static const char* usage_msg = "\ +Usage: %s NUM_INTERVALS [-r]\n\ +Spin on the GPU, logging intervals to stdout during which we are scheduled.\n\ + NUM_INTERVALS Number of intervals of execution to log (logs one less than\n\ + this number of preemptions).\n\ + -r, --raw Print raw logged GPU times (skip conversion to CPU time).\n"; + +int main(int argc, char **argv) { + cudaError_t err; // Needed for the SAFE() macro + struct interval *ivls_gpu, *ivls; + struct timespec start, end, end_ivls_only; + int num_ivls, skip_conversion; + uint64_t dev_ns = 0; + double d2h_scale = 1.0, host_s = 0; + pid_t pid = getpid(); + + // Argument parsing (could be improved) + if (argc != 2 && argc != 3) { + fprintf(stderr, usage_msg, argv[0]); + return 1; + } + num_ivls = atoi(argv[1]); + if (!num_ivls) { + fprintf(stderr, usage_msg, argv[0]); + return 1; + } + skip_conversion = argc == 3 && + (!strcmp(argv[2], "-r") || !strcmp(argv[2], "--raw")); + + // Synchronize the GPU and CPU clocks (if requested) using the utilities + // from task_host_utilities.cu. + if (!skip_conversion) { + // Note that this appears to almost always be 1.0 on recent GPUs + d2h_scale = GetGPUTimerScale(0); + // The following uses a subprocess for clock synchronization. Sketch? + GetHostDeviceTimeOffset(0, &host_s, &dev_ns); + } + + // Initialize GPU and CPU memory to store the intervals of execution + // Note that we use GPU memory rather than pinned host memory in order to make + // the time-monitoring loop on-GPU as fast as possible. (Pinned host memory is + // slower to read/write from/to.) + SAFE(cudaMalloc(&ivls_gpu, num_ivls * sizeof(struct interval))); + ivls = (struct interval*)malloc(num_ivls * sizeof(struct interval)); + if (!ivls) { + perror("While allocating memory"); + return 1; + } + + // Monitor execution intervals + fprintf(stderr, "(%d) Started monitoring execution for %d intervals...\n", + pid, num_ivls); + clock_gettime(CLOCK_MONOTONIC_RAW, &start); + loop_on_gpu<<<1,1>>>(ivls_gpu, num_ivls); + SAFE(cudaDeviceSynchronize()); + clock_gettime(CLOCK_MONOTONIC_RAW, &end_ivls_only); + + // Pull back results onto the CPU + SAFE(cudaMemcpy(ivls, ivls_gpu, num_ivls * sizeof(struct interval), + cudaMemcpyDeviceToHost)); + clock_gettime(CLOCK_MONOTONIC_RAW, &end); + + // Print out the intervals, the start and end times of our monitor, and some + // summary statistics + for (int i = 0; i < num_ivls; i++) { + if (skip_conversion) { + printf("%lu, %lu\n", ivls[i].start, ivls[i].end); + } else { + // Convert all the interval times into CPU time. + // First compute how far this interval's starting point is, in CPU seconds, + // from the point where the device and CPU clocks were synchronized + uint64_t start_time_since_sync = (ivls[i].start - dev_ns) * d2h_scale; + // Then convert into absolute CPU time via offset adjustment + uint64_t start = s2ns(host_s) + start_time_since_sync; + // Repeat the same process for interval end times + uint64_t end_time_since_sync = (ivls[i].end - dev_ns) * d2h_scale; + uint64_t end = s2ns(host_s) + end_time_since_sync; + printf("%lu, %lu\n", start, end); + } + } + fprintf(stderr, "(%d) Total monitor runtime: %ld ms\n", pid, + ns2ms(timediff(start, end))); + fprintf(stderr, "(%d) CPU logging start: %ld ns\n", pid, time2ns(start)); + fprintf(stderr, "(%d) CPU logging end: %ld ns\n", pid, time2ns(end)); + fprintf(stderr, "(%d) CPU copy out end: %ld ns\n", pid, time2ns(end)); + if (!skip_conversion) { + // Difference between CPU time of the start of the first interval and launch time + int64_t launch_oh = s2ns(host_s) + + (ivls[0].start - dev_ns) * d2h_scale - time2ns(start); + fprintf(stderr, "(%d) Aprox launch overhead: %ld ns\n", pid, launch_oh); + fprintf(stderr, "(%d) CPU clock - GPU clock: %ld tick gap\n", pid, + (long)s2ns(host_s) - dev_ns); + fprintf(stderr, "(%d) 1 CPU tick/1 GPU tick: %.3f\n", pid, d2h_scale); + } + + return 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 @@ +// 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 +#include +#include +#include +#include +#include +#include +#include +#include +#include + +// 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) { + // The following CUDA code is run in the child process + InternalReadGPUNanoseconds(cuda_device, shared_cpu_time, shared_gpu_time); + exit(0); + } + // The parent will wait for the child to finish, then return the value + // written to the shared buffer. + if (wait(&status) < 0) { + printf("Failed waiting on the child process.\n"); + FreeSharedBuffer(shared_cpu_time, sizeof(*shared_cpu_time)); + FreeSharedBuffer(shared_gpu_time, sizeof(*shared_gpu_time)); + return 0; + } + *host_seconds = *shared_cpu_time; + *gpu_nanoseconds = *shared_gpu_time; + FreeSharedBuffer(shared_cpu_time, sizeof(*shared_cpu_time)); + FreeSharedBuffer(shared_gpu_time, sizeof(*shared_gpu_time)); + if (!WIFEXITED(status)) { + printf("The child process didn't exit normally.\n"); + return 0; + } + return 1; +} + +// This function should always be run in a separate process. +static int InternalGetMaxResidentThreads(int cuda_device) { + struct cudaDeviceProp properties; + int warps_per_sm = 64; + if (!CheckCUDAError(cudaGetDeviceProperties(&properties, cuda_device))) { + return 0; + } + // Compute capability 2.0 devices have a 48 warps per SM. + if (properties.major <= 2) warps_per_sm = 48; + return warps_per_sm * properties.multiProcessorCount * properties.warpSize; +} + +int GetMaxResidentThreads(int cuda_device) { + int to_return, status; + pid_t pid = -1; + int *max_thread_count = NULL; + max_thread_count = (int *) AllocateSharedBuffer(sizeof(*max_thread_count)); + if (!max_thread_count) { + printf("Failed allocating shared buffer for IPC.\n"); + return 0; + } + pid = fork(); + if (pid < 0) { + printf("Failed creating a child process to get thread count: %s\n", + strerror(errno)); + return 0; + } + if (pid == 0) { + // The following CUDA code is run in the child process + *max_thread_count = InternalGetMaxResidentThreads(cuda_device); + exit(0); + } + // The parent will wait for the child to finish, then return the value + // written to the shared buffer. + if (wait(&status) < 0) { + printf("Failed waiting on the child process.\n"); + FreeSharedBuffer(max_thread_count, sizeof(*max_thread_count)); + return 0; + } + to_return = *max_thread_count; + FreeSharedBuffer(max_thread_count, sizeof(*max_thread_count)); + if (!WIFEXITED(status)) { + printf("The child process didn't exit normally.\n"); + return 0; + } + return to_return; +} + +#if __CUDA_ARCH__ >= 300 // Kepler+ +static __global__ void TimerSpin(uint64_t ns_to_spin) { + uint64_t start_time = GlobalTimer64(); + while ((GlobalTimer64() - start_time) < ns_to_spin) { + continue; + } +} +#elif __CUDA_ARCH__ < 200 // Tesla +// On sm_13, we /only/ have `clock` (32 bits) +static __device__ inline uint32_t Clock32(void) { + uint32_t lo_bits; + asm volatile("mov.u32 %0, %%clock;" : "=r"(lo_bits)); + return lo_bits; +} + +// 'clock' can easily roll over, so handle that for ancient architectures +static __global__ void TimerSpin(uint64_t ns_to_spin) { + uint64_t total_time = 0; + uint32_t last_time = Clock32(); + while (total_time < ns_to_spin) { + uint32_t time = Clock32(); + if (time < last_time) { + // Rollover. Compensate... + total_time += time; // rollover to now + total_time += UINT_MAX - last_time; // last to rollover + } else { + // Step counter + total_time += time - last_time; + } + last_time = time; + } +} +#else +// Could use clock64 for sm_2x (Fermi), but that's untested +#error Fermi-based GPUs (sm_2x) are unsupported! +#endif + +// This function is intended to be run in a child process. Returns -1 on error. +static double InternalGetGPUTimerScale(int cuda_device) { + struct timespec start, end; + uint64_t nanoseconds_elapsed; + if (!CheckCUDAError(cudaSetDevice(cuda_device))) return -1; + // Run the kernel once to warm up the GPU. + TimerSpin<<<1, 1>>>(1000); + if (!CheckCUDAError(cudaDeviceSynchronize())) return -1; + // After warming up, do the actual timing. + if (clock_gettime(CLOCK_MONOTONIC_RAW, &start) != 0) { + printf("Failed getting start time.\n"); + return -1; + } + TimerSpin<<<1, 1>>>(TIMER_SPIN_DURATION); + if (!CheckCUDAError(cudaDeviceSynchronize())) return -1; + if (clock_gettime(CLOCK_MONOTONIC_RAW, &end) != 0) { + printf("Failed getting end time.\n"); + return -1; + } + nanoseconds_elapsed = end.tv_sec * 1e9 + end.tv_nsec; + nanoseconds_elapsed -= start.tv_sec * 1e9 + start.tv_nsec; + return ((double) nanoseconds_elapsed) / ((double) TIMER_SPIN_DURATION); +} + +double GetGPUTimerScale(int cuda_device) { + double to_return; + double *scale = NULL; + int status; + pid_t pid; + scale = (double *) AllocateSharedBuffer(sizeof(*scale)); + if (!scale) { + printf("Failed allocating space to hold the GPU time scale.\n"); + return -1; + } + pid = fork(); + if (pid < 0) { + printf("Failed creating a child process.\n"); + FreeSharedBuffer(scale, sizeof(*scale)); + return -1; + } + if (pid == 0) { + // Access the GPU with the child process only. + *scale = InternalGetGPUTimerScale(cuda_device); + exit(0); + } + if (wait(&status) < 0) { + printf("Failed waiting on the child process.\n"); + FreeSharedBuffer(scale, sizeof(*scale)); + return -1; + } + to_return = *scale; + FreeSharedBuffer(scale, sizeof(*scale)); + if (!WIFEXITED(status)) { + printf("The child process didn't exit normally.\n"); + return -1; + } + return to_return; +} diff --git a/testbench.h b/testbench.h index 28d7f77..5b47bd4 100644 --- a/testbench.h +++ b/testbench.h @@ -11,4 +11,7 @@ #define s2ns(s) ((s)*1000l*1000l*1000l) #define ns2ms(s) ((s)/(1000l*1000l)) + +// Return the difference between two timestamps in nanoseconds #define timediff(start, end) ((s2ns((end).tv_sec) + (end).tv_nsec) - (s2ns((start).tv_sec) + (start).tv_nsec)) +#define time2ns(time) (s2ns((time).tv_sec) + (time).tv_nsec) -- cgit v1.2.2