mercredi 29 août 2018

Duplicate variadic template parameter

Context:

I am a Jr. Software Engineer, hopefully I am not reinventing the wheel, please let me know. I'd like to create a template function which wraps and calls another function element wise. For example:

// returns a*x + y
__device__ float saxpy(float a, float x, float y) {
  return a*x + y;
}


int main() {
  int A[4] = { 1,2,3,4 };
  int X[4] = { 1,2,3,4 };
  int Y[4] = { 1,1,1,1 };

  // A*X   = 1,4,9,16
  // A*X+Y = 2,5,10,17
  float *C = cudaReduce(saxpy, A, X, Y);

  for (int i = 0; i < 4; i++)
    printf("%d, ", C[i]); // should print "2, 5, 10, 17, "

  std::cin.ignore();
  return 0;
}

Importantly, I want to create this wrapper so that cuda calls are nicely wrapped when I perform element-wise operations. Though very incomplete, here is my pseudo-code attempt at the function wrapper.

I'd like to provide a minimal example; however, I have very little idea how to go about certain aspects of C++, so please forgive the large amounts of commented pseudocode:

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <iostream>

// returns a*x + y
__device__ float saxpy(float a, float x, float y) {
  return a*x + y;
}

// finds return type of function pointer
template<typename R, typename... A>
R ret(R(*)(A...));
template<typename C, typename R, typename... A>
R ret(R(C::*)(A...));

template<typename F, size_t N, typename... Args>
auto cudaReduce(F &f, Args(&...argsarray)[N]) {
  cudaSetDevice(0);

  // ret is function f's return type
  typedef decltype(ret(f)) ret;
  ret d_out[N], h_out[N];
  // cudaMalloc((void**)&d_out, sizeof(d_out));
  sendToCuda(argsarray...); // allocates and copies all contents of argsarray to cuda

  // reduceKernel<<<1, N>>>(f, d_out, dev_argsarray...);

  // cudaDeviceSynchronize();
  // cudaMemcpy(h_out, d_out, sizeof(h_out), cudaMemcpyDeviceToHost);
  // cudaFree(d_out);

  // for d_args in d_argsarray
  //   cudaFree(d_args);

  return h_out;
}

template<typename F, size_t N, typename Out, typename... Args>
__global__ void cudaReduceKernel(F &f, Out(&out)[N], Args(&...argsarray)[N]) {
  int tid = threadIdx.x;
  int i = tid + blockIdx.x * blockDim.x;

  // Below is invalid syntax; however, the 'pseudo-code' is what I'd like to achieve.
  // out[i] = f(argsarray[i]...);
}

// cuda malloc and memcpy
template<typename Arg, size_t N>
void sendToCuda(Arg(&args)[N]) {
  size_t buffer = sizeof(args);
  //cudaMalloc((void**)&dev_arg[ ??? ], buffer);
  //cudaMemcpy((void**)&dev_arg[ ??? ], args, buffer, cudaMemcpyHostToDevice);
}
template<typename Arg, size_t N, typename... Args>
void sendToCuda(Arg(&args)[N], Args(&...argsarray)[N]) {
  sendToCuda(args);
  sendToCuda(argsarray...);
}

int main() {
  int A[4] = { 1,2,3,4 };
  int X[4] = { 1,2,3,4 };
  int Y[4] = { 1,1,1,1 };

  // A*X   = 1,4,9,16
  // A*X+Y = 2,5,10,17
  float *C = cudaReduce(saxpy, A, X, Y);

  for (int i = 0; i < 4; i++)
    printf("%d, ", C[i]); // should print "2, 5, 10, 17, ", currently prints undefined behaviour

  std::cin.ignore();
  return 0;
}

I realize not everyone has time to completely review the code, so I will boil down the key problems into several points:

1. Is it possible to duplicate variadic template inputs, if so how? EX (not real code):

template<typename... Args>
void foo(Args... args) {
  Args2... args;
}

This is needed so that I can duplicate my input parameters to input parameters for my cuda malloc() and memcpy().

2. How would I go about the ith tuple of a variadic array parameter, like zipping in python. EX (not real code):

template<typename... Args, size_t N>
void bar(Args(&...argsarray)[N]) {
  // (python) ithvariadic = zip(*argsarray)[i]
  auto ithvariadic = argsarray[i]...;
}

Aucun commentaire:

Enregistrer un commentaire