#include <iostream>

using namespace std;

template<class T>
void quick_sort(T *data, size_t n) {
    if (n <= 1)return;

    size_t k = 1;
    for (size_t i = 1; i < n; i++) {
        if (data[i] < data[k - 1]) {
            swap(data[k - 1], data[i]);
            k++;
        }
    }

    quick_sort(data, k);
    quick_sort(data + k, n - k);
}

#include <llvm/IR/Module.h>
#include <llvm/IR/IRBuilder.h>
#include <llvm/IR/LLVMContext.h>
#include <llvm/Support/TargetSelect.h>
#include <llvm/Bitcode/BitcodeWriter.h>
#include <llvm/Support/raw_os_ostream.h>
#include <llvm/ExecutionEngine/ExecutionEngine.h>
#include <llvm/ExecutionEngine/GenericValue.h>

#include <string>
#include <vector>
#include <iostream>
#include <fstream>

using namespace std;
using namespace llvm;

int main() {
    LLVMContext context;
    IRBuilder<> builder(context);
    Module module("test", context);

    /*1.声明puts函数*/
    FunctionType *puts_type = FunctionType::get(builder.getInt32Ty(), {builder.getInt8PtrTy()}, true);
    Constant *puts_func = module.getOrInsertFunction("puts", puts_type);  // puts? printf?

    /*2.定义main函数*/
    FunctionType *main_type = FunctionType::get(builder.getInt32Ty(), false);
    Function *main_func = Function::Create(main_type, Function::ExternalLinkage, "main", &module);

    BasicBlock *main_body = BasicBlock::Create(context, "entry", main_func);
    builder.SetInsertPoint(main_body);

    CallInst *ret_value = builder.CreateCall(puts_func, builder.CreateGlobalStringPtr("Hello,world!!!"));
    builder.CreateRet(ret_value);

    /*3.output*/
    fstream fs("test.bc");  // lli test.bc
    raw_os_ostream out(fs);
    WriteBitcodeToFile(module, out);
    outs() << module;

    /*4.JIT*/
    InitializeNativeTarget();
    InitializeNativeTargetAsmPrinter();

    unique_ptr<Module> owner(&module);
    ExecutionEngine *engine = EngineBuilder(move(owner)).create();

    GenericValue gv = engine->runFunction(main_func, {});
    outs() << "Result: " << gv.IntVal;  // 获取函数调用结果

    return 0;
}

int a = 1024;
return a;

//申请栈空间，并返回对应的指针
AllocaInst *variable = builder.CreateAlloca(builder.getInt32Ty());
//将值存储到指定的位置
builder.CreateStore(builder.getInt32(1024), variable);
//加在指定位置处的值
builder.CreateRet(builder.CreateLoad(variable));

define i32 @main() {
entry:
  %0 = alloca i32
  store i32 1024, i32* %0
  %1 = load i32, i32* %0
  ret i32 %1
}

#include <iostream>

template<typename T>
T add(T a, T b) {
  std::cout << "T:";
  return a + b;
}

template<>
float add(float a, float b) {
  std::cout << "F:";
  return a + b; // 显示具体化
}

template
double add(double a, double b); // 显示实例化


int main() {
  std::cout << add(1, 2) << std::endl;      // 隐式实例化
  std::cout << add(1.f, 2.f) << std::endl;  // 显示具体化
  std::cout << add(1.0, 2.0) << std::endl;  // 显示实例化
  return 0;
}

// main.cu
#include <iostream>

#define CHECK(exp) \
if (exp != cudaSuccess) { \
printf("[ERROR] [%s:%d] %d\n", __FILE__, __LINE__, cudaGetLastError()); \
exit(-1); \
}

template<typename T>
__global__ void add(const T *a, const T *b, T *c, size_t n) {
  auto i = blockDim.x * blockIdx.x + threadIdx.x;
  if (i >= n)return;
  c[i] = a[i] + b[i];
}

int main() {
  float *a, *b, *c;
  size_t n = 1000;

  CHECK(cudaMallocHost(&a, n * sizeof(float)));
  CHECK(cudaMallocHost(&b, n * sizeof(float)));
  CHECK(cudaMallocHost(&c, n * sizeof(float)));

  for (int i = 0; i < n; i++) {
    a[i] = static_cast<float>(i);
    b[i] = static_cast<float>(i * i);
  }

  float *a_gpu, *b_gpu, *c_gpu;
  CHECK(cudaMalloc(&a_gpu, n * sizeof(float)));
  CHECK(cudaMalloc(&b_gpu, n * sizeof(float)));
  CHECK(cudaMalloc(&c_gpu, n * sizeof(float)));

  CHECK(cudaMemcpy(a_gpu, a, n * sizeof(float), cudaMemcpyHostToDevice));
  CHECK(cudaMemcpy(b_gpu, b, n * sizeof(float), cudaMemcpyHostToDevice));

  add<float><<<1, n>>>(a_gpu, b_gpu, c_gpu, n);

  CHECK(cudaMemcpy(c, c_gpu, n * sizeof(float), cudaMemcpyDeviceToHost));

  for (int i = 0; i < n; i++) {
    printf("%f ", c[i]);
  }
  printf("\n");

  cudaDeviceReset();
  return 0;
}

#ifndef CUDA_MAT_MAT_HPP
#define CUDA_MAT_MAT_HPP

#include <iostream>
#include <exception>
#include <sstream>

#include <cuda.h>
#include <cuda_runtime.h>

using namespace std;

#define CHECK(exp) \
if (exp != cudaSuccess) { \
    auto err = cudaGetLastError(); \
    printf("[ERROR] [%s:%d] %s(%d).\n", __FILE__, __LINE__, cudaGetErrorString(err), err); \
    exit(-1); \
}

template<class T>
class Mat {
  template<class K>
  friend ostream &operator<<(ostream &stream, const Mat<K> &mat);

public:
  Mat() : row_(0), col_(0), size_(0), cpu_data_(nullptr), gpu_data_(nullptr) {};

  Mat(size_t r, size_t c, bool use_cuda = false) : row_(r), col_(c), size_(r * c * sizeof(T)) {
    if (!use_cuda) {
      CHECK(cudaMallocHost(&cpu_data_, size_));
      gpu_data_ = nullptr;
    } else {
      cpu_data_ = nullptr;
      CHECK(cudaMalloc(&gpu_data_, size_));
    }
  }

  Mat(const initializer_list<T> &init_data) : Mat(1, init_data.size()) {
    T *q = cpu_data_;
    for (auto p = init_data.begin(); p != init_data.end(); p++, q++) {
      *q = *p;
    }
  }

  Mat(const initializer_list<initializer_list<T>> &init_data) : Mat(init_data.size(), init_data.begin()->size()) {
    T *q = cpu_data_;
    for (auto p1 = init_data.begin(); p1 != init_data.end(); p1++) {
      if (p1->size() != col_) {
        stringstream ss;
        ss << "The expected col is " << col_ << ", but the actual value is " << p1->size();
        throw invalid_argument(ss.str().c_str());
      }

      for (auto p2 = p1->begin(); p2 != p1->end(); p2++, q++) {
        *q = *p2;
      }
    }
  }

  Mat(const Mat &rhs) noexcept: row_(rhs.row_), col_(rhs.col_), size_(rhs.size_) {
    if (rhs.cpu_data_) {
      CHECK(cudaMallocHost(&cpu_data_, size_));
      CHECK(cudaMemcpy(cpu_data_, rhs.cpu_data_, size_, cudaMemcpyHostToHost));
    } else cpu_data_ = nullptr;

    if (rhs.gpu_data_) {
      CHECK(cudaMalloc(&gpu_data_, size_));
      CHECK(cudaMemcpy(gpu_data_, rhs.gpu_data_, size_, cudaMemcpyDeviceToDevice));
    } else gpu_data_ = nullptr;
  }

  Mat(Mat &&rhs) noexcept: row_(rhs.row_), col_(rhs.col_), size_(rhs.size_), cpu_data_(rhs.cpu_data_),
                           gpu_data_(rhs.gpu_data_) {
    rhs.col_ = rhs.row_ = rhs.row_ = 0;
    rhs.gpu_data_ = nullptr;
    rhs.cpu_data_ = nullptr;
  }

  Mat &operator=(const Mat &rhs) {
    Mat tmp(rhs);
    swap(tmp);
    return *this;
  }

  ~Mat() {
    if (cpu_data_)cudaFreeHost(cpu_data_);
    if (gpu_data_)cudaFree(gpu_data_);
  }

  Mat &cuda() {
    if (cpu_data_) {
      if (!gpu_data_)CHECK(cudaMalloc(&gpu_data_, size_));
      CHECK(cudaMemcpy(gpu_data_, cpu_data_, size_, cudaMemcpyHostToHost));
    }
    return *this;
  }

  Mat &cpu() {
    if (gpu_data_) {
      if (!cpu_data_)CHECK(cudaMallocHost(&cpu_data_, size_));
      CHECK(cudaMemcpy(cpu_data_, gpu_data_, size_, cudaMemcpyDeviceToHost));
    }
    return *this;
  }

  void swap(Mat<T> &rhs) {
    std::swap(row_, rhs.row_);
    std::swap(col_, rhs.col_);
    std::swap(size_, rhs.size_);
    std::swap(cpu_data_, rhs.cpu_data_);
    std::swap(gpu_data_, rhs.gpu_data_);
  }

  inline const T *cpu_data() const { return cpu_data_; }

  inline T *cpu_data_mutable() { return cpu_data_; }

  inline const T *gpu_data() const { return gpu_data_; };

  inline T *gpu_data_mutable() { return gpu_data_; }

  inline size_t row() const { return row_; }

  inline size_t col() const { return col_; }

private:
  T *cpu_data_, *gpu_data_;
  size_t row_, col_, size_;
};

template<class T>
ostream &operator<<(ostream &stream, const Mat<T> &mat) {
  if (mat.cpu_data_ == nullptr) {
    return stream;
  }

  T *p = mat.cpu_data_;
  for (size_t r = 0; r < mat.row_; r++) {
    for (size_t c = 0; c < mat.col_; c++) {
      stream << *p++ << " ";
    }
    stream << "\n";
  }

  return stream;
}
#endif //CUDA_MAT_MAT_HPP