path: root/task_host_utilities.cu

// Copyright 2017-2019 Nathan Otterness
// Copyright 2020-2024 Joshua Bakita
// This file was originally part of the cuda_scheduling_examiner tool, but has
// been copied and heavily modified for use in the gpu-microbench project.
#include <cuda_runtime.h>
#include <errno.h>
#include <sched.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);
}

// Note that CLOCK_MONOTONIC_RAW is slow before Linux 5.3 because it's not
// supported by vDSO.
static uint64_t CurrentNanoseconds(void) {
  struct timespec ts;
  if (clock_gettime(CLOCK_MONOTONIC_RAW, &ts) != 0) {
    printf("Error getting time.\n");
    exit(1);
  }
  return ((uint64_t) ts.tv_sec) * 1000*1000*1000 + ((uint64_t) ts.tv_nsec);
}

// 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;
  // The clock can roll over while we read the lo_bits in two cases:
  // 1. We read the hi_bits right as they changed, and lo_bits should be 0.
  // 2. We were preempted between reading hi_bits and hi_bits_2. In this case
  //    we need to re-read lo_bits, as they could be anything. Note that we
  //    assume that we'll never be suspended in an instruction-level preemption
  //    for more than ~4.29 seconds (which is how long it would take for the
  //    high bits to roll over again)---this allows us to just reread lo_bits.
  //    We also assume that this function will be preempted at at most one
  //    point during each invocation (the default timeslice length is 1-2 ms,
  //    so this assumption should always hold with the default scheduler).
  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 (hi_bits != hi_bits_2)
    asm volatile("mov.u32 %0, %%globaltimer_lo;" : "=r"(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 on CPU%d.\n",
    max_error * (1000.0 * 1000.0), sched_getcpu());
out:
  cudaFree(device_time);
  cudaFree((void*)start_barrier);
  cudaFree((void*)end_barrier);
  cudaFree((void*)ready_barrier);
}

// Returns 0 on success, 1 on error
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 __device__ 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 __device__ 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

// Waits on and sets barriers for the CPU before and after spinning for a
// specified number of clock ticks.
static __global__ void BarrierTimerSpin(uint64_t ns_to_spin, volatile uint32_t *ready_barrier,
    volatile uint32_t *start_barrier, volatile uint32_t *end_barrier) {