 }
 
 void ConcatCont(Tensor out, const std::vector<Tensor>& inputs, int axis) {
-  
+
 
   cudaSetDevice(out->getDevice().no);
   int step = 1;

  }
}

void TransposeND(Tensor out, Tensor in, const std::vector<int>& vAxis) {
  cudaSetDevice(out->getDevice().no);

  functional::Array<int, functional::Shape::size()> axes;
  int diff = functional::Shape::size() - vAxis.size();
  for(int i = 0; i < axes.size(); ++i)
    if(i < diff)
      axes[i] = i;
    else
      axes[i] = vAxis[i - diff] + diff;

  int length = out->shape().elements();
  int threads = std::min(MAX_THREADS, length);
  int blocks = std::min(MAX_BLOCKS, length / threads + (length % threads != 0));

  gTransposeND<<<blocks, threads>>>(out, in, axes);
}

__global__ void gSoftmax(float* out,

__global__ void gJoin2(float* out, size_t rowBatch, size_t cols,
                       const float* in1, size_t inStride1,
                       const float* in2, size_t inStride2) {

  int outStride = inStride1 + inStride2;
  int rows = rowBatch * outStride;

  for(int bid = 0; bid < rows; bid += gridDim.x) {
    int j = bid + blockIdx.x;
    if(j < rows) {

      float* rowOut = out + j * cols;

      int curBatch = j / outStride;
      int curPos = j % outStride;

