/*
GPU buffer transfer experiment code:
This code transfers chunks of data between buffers in main memory to test
the effects of bit-flips on energy consumption.

Copyright 2022 Texas State University

Redistribution and use in source and binary forms, with or without modification, are permitted provided that the following conditions are met:

1. Redistributions of source code must retain the above copyright notice, this list of conditions and the following disclaimer.

2. Redistributions in binary form must reproduce the above copyright notice, this list of conditions and the following disclaimer in the documentation and/or other materials provided with the distribution.

3. Neither the name of the copyright holder nor the names of its contributors may be used to endorse or promote products derived from this software without specific prior written permission.

THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.

Authors: Alex Fallin and Martin Burtscher

URL: The latest version of this code is available at
https://cs.txstate.edu/~burtscher/research/bit-flips/.
*/

#include <cstdio>
#include <algorithm>
#include <cuda.h>
#include <sys/time.h>
#include <stdlib.h>
#include <unistd.h>
#include <tgmath.h>
#include "gpu_energy_monitor.h"

static const int ThreadsPerBlock = 256;
static const double RUN_TIME = 10;
static const int TOT_RUNS = 10;

// source of hash function: https://stackoverflow.com/questions/664014/what-integer-hash-function-are-good-that-accepts-an-integer-hash-key
static __device__ unsigned int hash(unsigned int val) {
  val = ((val >> 16) ^ val) * 0x45d9f3b;
  val = ((val >> 16) ^ val) * 0x45d9f3b;
  return (val >> 16) ^ val;
}

// Fills the buffer with a buff_sz-word aligned pattern that causes 100% bitflips
static __global__ void init_eq_oz_block(const int size, int* const buf, const int buf_sz) {
  const int i = threadIdx.x + blockIdx.x * ThreadsPerBlock;
  if (i < size) {
    if (i % (buf_sz * 2) < buf_sz) {
      buf[i] = 0xFFFFFFFF;
    } else {
      buf[i] = 0x00000000;
    }
  }
}

// Fills the buffer with an equal number of 1s and zeros that do not cause flips
static __global__ void init_eq_oz_inter(const int size, int* const buf) {
  const int i = threadIdx.x + blockIdx.x * ThreadsPerBlock;
  if (i < size) {
    buf[i] = 0xAAAAAAAA;
  }
}

// Fills the buffer with all 0s
static __global__ void init_all_0(const int size, int* const buf) {
  const int i = threadIdx.x + blockIdx.x * ThreadsPerBlock;
  if (i < size) {
    buf[i] = 0x00000000;
  }
}

// Fills the buffer with all 1s
static __global__ void init_all_1(const int size, int* const buf) {
  const int i = threadIdx.x + blockIdx.x * ThreadsPerBlock;
  if (i < size) {
    buf[i] = 0xFFFFFFFF;
  }
}

// Fills the buffer with random data
static __global__ void init_random(const int size, int* const buf) {
  const int i = threadIdx.x + blockIdx.x * ThreadsPerBlock;
  if (i < size) {
    buf[i] = hash(i + 1);  // random numbers
  }
}

static __global__ void copy(const int size, const int* const buf1, int* const buf2) {
  const int i = threadIdx.x + blockIdx.x * ThreadsPerBlock;
  if (i < size) {
    buf2[i] = buf1[i];
  }
}

static void CheckCuda() {
  cudaError_t e;
  cudaDeviceSynchronize();
  if (cudaSuccess != (e = cudaGetLastError())) {
    fprintf(stderr, "CUDA error %d: %s\n", e, cudaGetErrorString(e));
    exit(-1);
  }
}

// Copies the buffer back and forth until RUNTIME seconds have passed
static int copy_rep(int size, int* d_buf1, int* d_buf2, int rep) {
  timeval start, end;
  gettimeofday(&start, NULL);
  gettimeofday(&end, NULL);

  if (rep == -1) {
    int num_reps = 0;
    while ((end.tv_sec - start.tv_sec + (end.tv_usec - start.tv_usec) / 1000000.0) < RUN_TIME) {
      copy<<<(size + ThreadsPerBlock - 1) / ThreadsPerBlock, ThreadsPerBlock>>>(size, d_buf1, d_buf2);
      cudaDeviceSynchronize();
      std::swap(d_buf1, d_buf2);
      gettimeofday(&end, NULL);
      num_reps++;
    }
    return num_reps;
  } else {
    for (int i = 0; i < rep; i++) {
      copy<<<(size + ThreadsPerBlock - 1) / ThreadsPerBlock, ThreadsPerBlock>>>(size, d_buf1, d_buf2);
      cudaDeviceSynchronize();
      std::swap(d_buf1, d_buf2);
    }
    return 0;
  }
}

// An optional function that warms the GPU to a target temperature before progressing. This improves result stability
static void warm_up_gpu(int size, int* d_buf1, int* d_buf2, GPUMonitor gpu, const int target_temp) {
  // Running the GPU to get it up to temp
  unsigned int curr_temp = 0;
  init_random<<<(size + ThreadsPerBlock - 1) / ThreadsPerBlock, ThreadsPerBlock>>>(size, d_buf1);

  while (curr_temp <= target_temp) {
    copy_rep(size, d_buf1, d_buf2, -1);
    curr_temp = gpu.getTemperature();
    // To prevent the temperature readings from messing up the first energy reading
    cudaDeviceSynchronize();
  }
}

int main(int argc, char* argv[]) {
  printf("Best case vs worst case consumption\n");

  if (argc != 4) {
    fprintf(stderr, "USAGE: %s gpu_num size_in_bytes warmup_temp\n", argv[0]);
    exit(-1);
  }

  int cuda_device = atoi(argv[1]);
  cudaSetDevice(cuda_device);
  cudaDeviceProp deviceProp;
  cudaGetDeviceProperties(&deviceProp, cuda_device);
  if ((deviceProp.major == 9999) && (deviceProp.minor == 9999)) {
    printf("ERROR: there is no CUDA capable device\n\n");
    exit(-1);
  }
  const int SMs = deviceProp.multiProcessorCount;
  const int mTpSM = deviceProp.maxThreadsPerMultiProcessor;
  printf("gpu: %s with %d SMs and %d mTpSM (%.1f MHz and %.1f MHz)\n", deviceProp.name, SMs, mTpSM, deviceProp.clockRate * 0.001, deviceProp.memoryClockRate * 0.001);

  int size = atoi(argv[2]);
  int warmup_temp = atoi(argv[3]);

  int* d_buf1;
  if (cudaSuccess != cudaMalloc((void**) &d_buf1, size * sizeof(int))) {
    fprintf(stderr, "ERROR: could not allocate memory\n");
    exit(-1);
  }
  int* d_buf2;
  if (cudaSuccess != cudaMalloc((void**) &d_buf2, size * sizeof(int))) {
    fprintf(stderr, "ERROR: could not allocate memory\n");
    exit(-1);
  }

  GPUMonitor gpu;
  gpu.initMonitor(cuda_device);
  timeval start, end;

  warm_up_gpu(size, d_buf1, d_buf2, gpu, warmup_temp);

  // Get num reps
  init_random<<<(size + ThreadsPerBlock - 1) / ThreadsPerBlock, ThreadsPerBlock>>>(size, d_buf1);
  int reps = copy_rep(size, d_buf1, d_buf2, -1);

  double rand_runtimes[TOT_RUNS];
  unsigned long long rand_cons[TOT_RUNS];
  double zeros_runtimes[TOT_RUNS];
  unsigned long long zeros_cons[TOT_RUNS];
  double ones_runtimes[TOT_RUNS];
  unsigned long long ones_cons[TOT_RUNS];

  for (int r = 0; r < TOT_RUNS; r++) {
    /*
    * Best-case 1s
    */
    init_all_1<<<(size + ThreadsPerBlock - 1) / ThreadsPerBlock, ThreadsPerBlock>>>(size, d_buf1);

    gpu.startEnergy();

    gettimeofday(&start, NULL);

    copy_rep(size, d_buf1, d_buf2, reps);

    gettimeofday(&end, NULL);
    ones_runtimes[r] = end.tv_sec - start.tv_sec + (end.tv_usec - start.tv_usec) / 1000000.0;

    ones_cons[r] = gpu.getEnergyConsumption();

    /*
    * Best-case 0s
    */
    init_all_0<<<(size + ThreadsPerBlock - 1) / ThreadsPerBlock, ThreadsPerBlock>>>(size, d_buf1);

    gpu.startEnergy();

    gettimeofday(&start, NULL);

    copy_rep(size, d_buf1, d_buf2, reps);

    gettimeofday(&end, NULL);
    zeros_runtimes[r] = end.tv_sec - start.tv_sec + (end.tv_usec - start.tv_usec) / 1000000.0;

    zeros_cons[r] = gpu.getEnergyConsumption();

    /*
    * Random
    */
    init_random<<<(size + ThreadsPerBlock - 1) / ThreadsPerBlock, ThreadsPerBlock>>>(size, d_buf1);

    gpu.startEnergy();

    gettimeofday(&start, NULL);

    copy_rep(size, d_buf1, d_buf2, reps);

    gettimeofday(&end, NULL);
    rand_runtimes[r] = end.tv_sec - start.tv_sec + (end.tv_usec - start.tv_usec) / 1000000.0;

    rand_cons[r] = gpu.getEnergyConsumption();
  }

  // Print results
  // reps, rand_con, rand_time, 1s_con, 1s_time, 0s_con, 0s_time
  for (int run = 0; run < TOT_RUNS; run++) {
    printf("%d, %lld, %.4f, %lld, %.4f, %lld, %.4f\n", reps, rand_cons[run], rand_runtimes[run], ones_cons[run], ones_runtimes[run], zeros_cons[run], zeros_runtimes[run]);
  }

  CheckCuda();

  cudaFree(d_buf1);
  cudaFree(d_buf2);

  return 0;
}