#include "tensoroperations.cuh"

const int THREAD_SIZE_XY = 1 << 10;
const int THREAD_SIZE_Z = 1 << 6;

__global__
void addElementwiseD(int size, float* ptr1, float* ptr2, float* ptr3) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < size) ptr3[idx] = ptr1[idx] + ptr2[idx];
}

std::unique_ptr<float[]> CUDAaddElementwise(std::unique_ptr<float[]>& in_ptr1, std::unique_ptr<float[]>& in_ptr2, int ptr_size) {
    int gpu_ptr1_bytes = ptr_size * sizeof(float);

    float* gpu_ptr1;
    float* gpu_ptr2;
    float* gpu_ptr3;
    cudaMalloc(&gpu_ptr1, gpu_ptr1_bytes);
    cudaMalloc(&gpu_ptr2, gpu_ptr1_bytes);
    cudaMalloc(&gpu_ptr3, gpu_ptr1_bytes);

    cudaMemcpy(gpu_ptr1, in_ptr1.get(), gpu_ptr1_bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(gpu_ptr2, in_ptr2.get(), gpu_ptr1_bytes, cudaMemcpyHostToDevice);

    int dimGridX = (ptr_size + THREAD_SIZE_XY - 1) / THREAD_SIZE_XY;
    addElementwiseD <<< dimGridX, THREAD_SIZE_XY >>> (ptr_size, gpu_ptr1, gpu_ptr2, gpu_ptr3);

    std::unique_ptr<float[]> out_ptr(new float[ptr_size]);
    cudaMemcpy(out_ptr.get(), gpu_ptr3, gpu_ptr1_bytes, cudaMemcpyDeviceToHost);

    cudaFree(gpu_ptr1);
    cudaFree(gpu_ptr2);
    cudaFree(gpu_ptr3);
   
    return out_ptr;
}

__global__
void subtractElementwiseD(int size, float* ptr1, float* ptr2, float* ptr3) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < size) ptr3[idx] = ptr1[idx] - ptr2[idx];
}

std::unique_ptr<float[]> CUDAsubtractElementwise(std::unique_ptr<float[]>& in_ptr1, std::unique_ptr<float[]>& in_ptr2, int ptr_size) {
    int gpu_ptr1_bytes = ptr_size * sizeof(float);

    float* gpu_ptr1;
    float* gpu_ptr2;
    float* gpu_ptr3;
    cudaMalloc(&gpu_ptr1, gpu_ptr1_bytes);
    cudaMalloc(&gpu_ptr2, gpu_ptr1_bytes);
    cudaMalloc(&gpu_ptr3, gpu_ptr1_bytes);

    cudaMemcpy(gpu_ptr1, in_ptr1.get(), gpu_ptr1_bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(gpu_ptr2, in_ptr2.get(), gpu_ptr1_bytes, cudaMemcpyHostToDevice);

    int dimGridX = (ptr_size + THREAD_SIZE_XY - 1) / THREAD_SIZE_XY;
    subtractElementwiseD <<< dimGridX, THREAD_SIZE_XY >>> (ptr_size, gpu_ptr1, gpu_ptr2, gpu_ptr3);

    std::unique_ptr<float[]> out_ptr(new float[ptr_size]);
    cudaMemcpy(out_ptr.get(), gpu_ptr3, gpu_ptr1_bytes, cudaMemcpyDeviceToHost);

    cudaFree(gpu_ptr1);
    cudaFree(gpu_ptr2);
    cudaFree(gpu_ptr3);
   
    return out_ptr;
}

__global__
void multiplyElementwiseD(int size, float* ptr1, float* ptr2, float* ptr3) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < size) ptr3[idx] = ptr1[idx] * ptr2[idx];
}

std::unique_ptr<float[]> CUDAmultiplyElementwise(std::unique_ptr<float[]>& in_ptr1, std::unique_ptr<float[]>& in_ptr2, int ptr_size) {
    int gpu_ptr1_bytes = ptr_size * sizeof(float);

    float* gpu_ptr1;
    float* gpu_ptr2;
    float* gpu_ptr3;
    cudaMalloc(&gpu_ptr1, gpu_ptr1_bytes);
    cudaMalloc(&gpu_ptr2, gpu_ptr1_bytes);
    cudaMalloc(&gpu_ptr3, gpu_ptr1_bytes);

    cudaMemcpy(gpu_ptr1, in_ptr1.get(), gpu_ptr1_bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(gpu_ptr2, in_ptr2.get(), gpu_ptr1_bytes, cudaMemcpyHostToDevice);

    int dimGridX = (ptr_size + THREAD_SIZE_XY - 1) / THREAD_SIZE_XY;
    multiplyElementwiseD <<< dimGridX, THREAD_SIZE_XY >>> (ptr_size, gpu_ptr1, gpu_ptr2, gpu_ptr3);

    std::unique_ptr<float[]> out_ptr(new float[ptr_size]);
    cudaMemcpy(out_ptr.get(), gpu_ptr3, gpu_ptr1_bytes, cudaMemcpyDeviceToHost);

    cudaFree(gpu_ptr1);
    cudaFree(gpu_ptr2);
    cudaFree(gpu_ptr3);
   
    return out_ptr;
}

__global__
void divideElementwiseD(int size, float* ptr1, float* ptr2, float* ptr3) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < size) ptr3[idx] = ptr1[idx] / ptr2[idx];
}

std::unique_ptr<float[]> CUDAdivideElementwise(std::unique_ptr<float[]>& in_ptr1, std::unique_ptr<float[]>& in_ptr2, int ptr_size) {
    int gpu_ptr1_bytes = ptr_size * sizeof(float);

    float* gpu_ptr1;
    float* gpu_ptr2;
    float* gpu_ptr3;
    cudaMalloc(&gpu_ptr1, gpu_ptr1_bytes);
    cudaMalloc(&gpu_ptr2, gpu_ptr1_bytes);
    cudaMalloc(&gpu_ptr3, gpu_ptr1_bytes);

    cudaMemcpy(gpu_ptr1, in_ptr1.get(), gpu_ptr1_bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(gpu_ptr2, in_ptr2.get(), gpu_ptr1_bytes, cudaMemcpyHostToDevice);

    int dimGridX = (ptr_size + THREAD_SIZE_XY - 1) / THREAD_SIZE_XY;
    divideElementwiseD <<< dimGridX, THREAD_SIZE_XY >>> (ptr_size, gpu_ptr1, gpu_ptr2, gpu_ptr3);

    std::unique_ptr<float[]> out_ptr(new float[ptr_size]);
    cudaMemcpy(out_ptr.get(), gpu_ptr3, gpu_ptr1_bytes, cudaMemcpyDeviceToHost);

    cudaFree(gpu_ptr1);
    cudaFree(gpu_ptr2);
    cudaFree(gpu_ptr3);
   
    return out_ptr;
}

__global__
void powerElementwiseD(int size, float* ptr1, float* ptr2, float* ptr3) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < size) ptr3[idx] = std::pow(ptr1[idx], ptr2[idx]);
}

std::unique_ptr<float[]> CUDApowerElementwise(std::unique_ptr<float[]>& in_ptr1, std::unique_ptr<float[]>& in_ptr2, int ptr_size) {
    int gpu_ptr1_bytes = ptr_size * sizeof(float);

    float* gpu_ptr1;
    float* gpu_ptr2;
    float* gpu_ptr3;
    cudaMalloc(&gpu_ptr1, gpu_ptr1_bytes);
    cudaMalloc(&gpu_ptr2, gpu_ptr1_bytes);
    cudaMalloc(&gpu_ptr3, gpu_ptr1_bytes);

    cudaMemcpy(gpu_ptr1, in_ptr1.get(), gpu_ptr1_bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(gpu_ptr2, in_ptr2.get(), gpu_ptr1_bytes, cudaMemcpyHostToDevice);

    int dimGridX = (ptr_size + THREAD_SIZE_XY - 1) / THREAD_SIZE_XY;
    powerElementwiseD <<< dimGridX, THREAD_SIZE_XY >>> (ptr_size, gpu_ptr1, gpu_ptr2, gpu_ptr3);

    std::unique_ptr<float[]> out_ptr(new float[ptr_size]);
    cudaMemcpy(out_ptr.get(), gpu_ptr3, gpu_ptr1_bytes, cudaMemcpyDeviceToHost);

    cudaFree(gpu_ptr1);
    cudaFree(gpu_ptr2);
    cudaFree(gpu_ptr3);
   
    return out_ptr;
}

__global__
void transposeD(int cols, int rows, int depths, float* ptr1, float* ptr2) {
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int depth = blockIdx.z * blockDim.z + threadIdx.z;
    // Of course this is going to need a Z coordinate for the infinite dimensions it can take
    if ((col < cols) && (row < rows) && (depth < depths)) ptr2[depth * rows * cols + row * cols + col] = ptr1[depth * rows * cols + col * rows + row];
}

// I need to reformat all of the other functions to fit this
std::unique_ptr<float[]> CUDAtranspose(std::unique_ptr<float[]>& in_ptr1, std::unique_ptr<int[]>& in_ptr1_dims, int in_ptr1_dims_size, int ptr1_size) {
    int cols = in_ptr1_dims[0];
    int rows = in_ptr1_dims[1];
    // Is there a faster way to do this
    int depths = 1;
    for (int i = 2; i < in_ptr1_dims_size; i++) {
        depths *= in_ptr1_dims[i];
    }

    int gpu_ptr1_bytes = ptr1_size * sizeof(float);

    float* gpu_ptr1;
    float* gpu_ptr2;
    cudaMalloc(&gpu_ptr1, gpu_ptr1_bytes);
    cudaMalloc(&gpu_ptr2, gpu_ptr1_bytes);
    cudaMemcpy(gpu_ptr1, in_ptr1.get(), gpu_ptr1_bytes, cudaMemcpyHostToDevice);

    int grid_cols = (cols + std::sqrt(THREAD_SIZE_XY / THREAD_SIZE_Z) - 1) / std::sqrt(THREAD_SIZE_XY / THREAD_SIZE_Z);
    int grid_rows = (rows + std::sqrt(THREAD_SIZE_XY / THREAD_SIZE_Z) - 1) / std::sqrt(THREAD_SIZE_XY / THREAD_SIZE_Z);
    int grid_depths = (depths + THREAD_SIZE_Z - 1) / THREAD_SIZE_Z;

    dim3 gridSize(grid_cols, grid_cols, grid_depths);
    dim3 threadSize(std::sqrt(THREAD_SIZE_XY / THREAD_SIZE_Z), std::sqrt(THREAD_SIZE_XY / THREAD_SIZE_Z), THREAD_SIZE_Z);

    transposeD <<< gridSize, threadSize >>> (cols, rows, depths, gpu_ptr1, gpu_ptr2);

    std::unique_ptr<float[]> out_ptr(new float[ptr1_size]);
    cudaMemcpy(out_ptr.get(), gpu_ptr2, gpu_ptr1_bytes, cudaMemcpyDeviceToHost);

    cudaFree(gpu_ptr1);
    cudaFree(gpu_ptr2);
   
    return out_ptr;
} 

__global__
void multiplyD(int cols, int shared, int rows, int depths, float* ptr1, float* ptr2, float* ptr3) {
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int depth = blockIdx.z * blockDim.z + threadIdx.z;

    float sum;
    if ((col < cols) && (row < rows) && (depth < depths)) {
        sum = 0;
        for (int i = 0; i < shared; i++) {
            sum += ptr1[depth * rows * cols + row * shared + i] * ptr2[depth * rows * cols + i * cols + col];
        }
        ptr3[depth * rows * cols + row * cols + col] = sum;
    }
}

std::unique_ptr<float[]> CUDAmultiply(std::unique_ptr<float[]>& in_ptr1, std::unique_ptr<int[]>& in_ptr1_dims, int in_ptr1_dims_size, int ptr1_size, std::unique_ptr<float[]>& in_ptr2, std::unique_ptr<int[]>& in_ptr2_dims, int in_ptr2_dims_size, int ptr2_size) {
    int ptr1_rows = in_ptr1_dims[1];
    int ptr2_cols = in_ptr2_dims[0];
    int shared_size = in_ptr1_dims[0];
    int depths = 1;
    for (int i = 2; i < in_ptr1_dims_size; i++) { // In theory the rest of the dims after should be the exact same if we assume they are correct
        depths *= in_ptr1_dims[i];
    } 

    int gpu_ptr1_bytes = ptr1_size * sizeof(float);
    int gpu_ptr2_bytes = ptr2_size * sizeof(float);
    int gpu_ptr3_bytes = depths * ptr1_rows * ptr2_cols * sizeof(float);

    float* gpu_ptr1;
    float* gpu_ptr2;
    float* gpu_ptr3;
    cudaMalloc(&gpu_ptr1, gpu_ptr1_bytes);
    cudaMalloc(&gpu_ptr2, gpu_ptr2_bytes);
    cudaMalloc(&gpu_ptr3, gpu_ptr3_bytes);
    cudaMemcpy(gpu_ptr1, in_ptr1.get(), gpu_ptr1_bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(gpu_ptr2, in_ptr2.get(), gpu_ptr2_bytes, cudaMemcpyHostToDevice);

    int grid_cols = (ptr2_cols + std::sqrt(THREAD_SIZE_XY / THREAD_SIZE_Z) - 1) / std::sqrt(THREAD_SIZE_XY / THREAD_SIZE_Z);
    int grid_rows = (ptr1_rows + std::sqrt(THREAD_SIZE_XY / THREAD_SIZE_Z) - 1) / std::sqrt(THREAD_SIZE_XY / THREAD_SIZE_Z);
    int grid_depths = (depths + THREAD_SIZE_Z - 1) / THREAD_SIZE_Z;

    dim3 gridSize(grid_cols, grid_cols, grid_depths);
    dim3 threadSize(std::sqrt(THREAD_SIZE_XY / THREAD_SIZE_Z), std::sqrt(THREAD_SIZE_XY / THREAD_SIZE_Z), THREAD_SIZE_Z);

    multiplyD <<< gridSize, threadSize >>> (ptr2_cols, shared_size, ptr1_rows, depths, gpu_ptr1, gpu_ptr2, gpu_ptr3);

    std::unique_ptr<float[]> out_ptr(new float[depths * ptr1_rows * ptr2_cols]);
    cudaMemcpy(out_ptr.get(), gpu_ptr3, gpu_ptr3_bytes, cudaMemcpyDeviceToHost);

    cudaFree(gpu_ptr1);
    cudaFree(gpu_ptr2);
    cudaFree(gpu_ptr3);

    return out_ptr;
}

__global__
void rotateD(int cols, int rows, int depths, float* ptr1, float* ptr2) {
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int depth = blockIdx.z * blockDim.z + threadIdx.z;

    if ((col < cols) && (row < rows) && (depth < depths)) ptr2[depth * rows * cols + (rows - row - 1) * cols + (cols - col - 1)] = ptr1[depth * rows * cols + row * cols + col];
}

std::unique_ptr<float[]> CUDArotate(std::unique_ptr<float[]>& in_ptr1, std::unique_ptr<int[]>& in_ptr1_dims, int in_ptr1_dims_size, int ptr1_size) {
    int ptr1_cols = in_ptr1_dims[0];
    int ptr1_rows = in_ptr1_dims[1];
    int depths = 1;
    for (int i = 2; i < in_ptr1_dims_size; i++) {
        depths *= in_ptr1_dims[i];
    }

    int gpu_ptr_bytes = ptr1_size * sizeof(float);

    float* gpu_ptr1;
    float* gpu_ptr2;
    cudaMalloc(&gpu_ptr1, gpu_ptr_bytes);
    cudaMalloc(&gpu_ptr2, gpu_ptr_bytes);
    cudaMemcpy(gpu_ptr1, in_ptr1.get(), gpu_ptr_bytes, cudaMemcpyHostToDevice);

    int grid_cols = (ptr1_cols + std::sqrt(THREAD_SIZE_XY / THREAD_SIZE_Z) - 1) / std::sqrt(THREAD_SIZE_XY / THREAD_SIZE_Z);
    int grid_rows = (ptr1_rows + std::sqrt(THREAD_SIZE_XY / THREAD_SIZE_Z) - 1) / std::sqrt(THREAD_SIZE_XY / THREAD_SIZE_Z);
    int grid_depths = (depths + THREAD_SIZE_Z - 1) / THREAD_SIZE_Z;

    dim3 gridSize(grid_cols, grid_cols, grid_depths);
    dim3 threadSize(std::sqrt(THREAD_SIZE_XY / THREAD_SIZE_Z), std::sqrt(THREAD_SIZE_XY / THREAD_SIZE_Z), THREAD_SIZE_Z);

    rotateD <<< gridSize, threadSize >>> (ptr1_cols, ptr1_rows, depths, gpu_ptr1, gpu_ptr2);

    std::unique_ptr<float[]> out_ptr(new float[ptr1_size]);
    cudaMemcpy(out_ptr.get(), gpu_ptr2, gpu_ptr_bytes, cudaMemcpyDeviceToHost);

    cudaFree(gpu_ptr1);
    cudaFree(gpu_ptr2);

    return out_ptr;
}

__global__
void maxPoolingD(int cols, int rows, int depths, int kernel_cols, int kernel_rows, int stride_cols, int stride_rows, float* ptr1, float* ptr2) {
    int col = blockIdx.x * blockDim.x + threadIdx.x; // Col of the unpooled ptr
    int row = blockIdx.y * blockDim.y + threadIdx.y; // Row of the unpooled ptr
    int depth = blockIdx.z * blockDim.z + threadIdx.z; // Depth of the unpooled ptr

    if ((col < cols - kernel_cols + 1) && (row < rows - kernel_rows + 1) && (depth < depths)) {
        if ((col % stride_cols == 0) && (row % stride_rows == 0)) {

            int max = ptr1[depth * rows * cols + row * cols + col];
            int comparison;

            for (int i = 0; i < kernel_rows; i++) {
                for (int j = 0; j < kernel_cols; j++) {

                    comparison = ptr1[depth * rows * cols + (row + i) * cols + (col + j)];
                    if (max < comparison) max = comparison; 
                    
                }
            }

            int pooled_cols_size = (cols - kernel_cols + stride_cols) / stride_cols;
            int pooled_rows_size = (rows - kernel_rows + stride_rows) / stride_rows;

            int pooled_col = (col - kernel_cols + stride_cols) / stride_cols;
            if (pooled_col < 0) pooled_col = 0;
            int pooled_row = (row - kernel_rows + stride_rows) / stride_rows;
            if (pooled_row < 0) pooled_row = 0;

            ptr2[depth * pooled_rows_size * pooled_cols_size + pooled_row * pooled_cols_size + pooled_col] = max;
        }
    }
}

std::unique_ptr<float[]> CUDAmaxPooling(std::unique_ptr<float[]>& in_ptr1, std::unique_ptr<int[]>& in_ptr1_dims, int in_ptr1_dims_size, int ptr1_size, int kernel_cols, int kernel_rows, int stride_cols, int stride_rows) {
    // This is raw dimensions
    int ptr1_cols = in_ptr1_dims[0];
    int ptr1_rows = in_ptr1_dims[1];
    int depths = 1;
    for (int i = 2; i < in_ptr1_dims_size; i++) {
        depths *= in_ptr1_dims[i];
    }

    // This is pooled dims
    int ptr2_cols = (ptr1_cols - kernel_cols + stride_cols) / stride_cols;
    int ptr2_rows = (ptr1_rows - kernel_rows + stride_rows) / stride_rows;
    int ptr2_size = ptr2_cols * ptr2_rows * depths;

    int gpu_ptr1_bytes = ptr1_size * sizeof(float);
    int gpu_ptr2_bytes = ptr2_size * sizeof(float);

    float* gpu_ptr1;
    float* gpu_ptr2;
    cudaMalloc(&gpu_ptr1, gpu_ptr1_bytes);
    cudaMalloc(&gpu_ptr2, gpu_ptr2_bytes);
    cudaMemcpy(gpu_ptr1, in_ptr1.get(), gpu_ptr1_bytes, cudaMemcpyHostToDevice);

    int grid_cols = (ptr1_cols + std::sqrt(THREAD_SIZE_XY / THREAD_SIZE_Z) - 1) / std::sqrt(THREAD_SIZE_XY / THREAD_SIZE_Z);
    int grid_rows = (ptr1_rows + std::sqrt(THREAD_SIZE_XY / THREAD_SIZE_Z) - 1) / std::sqrt(THREAD_SIZE_XY / THREAD_SIZE_Z);
    int grid_depths = (depths + THREAD_SIZE_Z - 1) / THREAD_SIZE_Z;

    dim3 gridSize(grid_cols, grid_cols, grid_depths);
    dim3 threadSize(std::sqrt(THREAD_SIZE_XY / THREAD_SIZE_Z), std::sqrt(THREAD_SIZE_XY / THREAD_SIZE_Z), THREAD_SIZE_Z);

    maxPoolingD <<< gridSize, threadSize >>> (ptr1_cols, ptr1_rows, depths, kernel_rows, kernel_cols, stride_cols, stride_cols, gpu_ptr1, gpu_ptr2);

    std::unique_ptr<float[]> out_ptr(new float[ptr2_size]);
    cudaMemcpy(out_ptr.get(), gpu_ptr2, gpu_ptr2_bytes, cudaMemcpyDeviceToHost);

    cudaFree(gpu_ptr1);
    cudaFree(gpu_ptr2);
    
    return out_ptr;
}

__global__
void poolingDerivD(int cols, int rows, int depths, int kernel_cols, int kernel_rows, int stride_cols, int stride_rows, float* ptr1, float* ptr2, float* ptr3) {
    int col = blockIdx.x * blockDim.x + threadIdx.x; // Col of the unpooled ptr
    int row = blockIdx.y * blockDim.y + threadIdx.y; // Row of the unpooled ptr
    int depth = blockIdx.z * blockDim.z + threadIdx.z; // Depth of the unpooled ptr

    if ((col < cols - kernel_cols + 1) && (row < rows - kernel_rows + 1) && (depth < depths)) {
        if ((col % stride_cols == 0) && (row % stride_rows == 0)) {

            int max = ptr1[depth * rows * cols + row * cols + col];
            int argmax_col = 0;
            int argmax_row = 0;
            int comparison;

            for (int i = 0; i < kernel_rows; i++) {
                for (int j = 0; j < kernel_cols; j++) {

                    comparison = ptr1[depth * rows * cols + (row + i) * cols + (col + j)];
                    if (max < comparison) {
                        max = comparison;
                        argmax_col = j;
                        argmax_row = i;
                    }

                }
            } 

            int pooled_cols_size = (cols - kernel_cols + stride_cols) / stride_cols;
            int pooled_rows_size = (rows - kernel_rows + stride_rows) / stride_rows;

            int pooled_col = (col - kernel_cols + stride_cols) / stride_cols;
            if (pooled_col < 0) pooled_col = 0;
            int pooled_row = (row - kernel_rows + stride_rows) / stride_rows;
            if (pooled_row < 0) pooled_row = 0;

            ptr3[depth * rows * cols + (row + argmax_row) * cols + (col + argmax_col)] += ptr2[depth * pooled_rows_size * pooled_cols_size + pooled_row * pooled_cols_size + pooled_col];

        }
    }
}

std::unique_ptr<float[]> CUDApoolingDeriv(std::unique_ptr<float[]>& in_ptr1, std::unique_ptr<int[]>& in_ptr1_dims, int in_ptr1_dims_size, int ptr1_size, std::unique_ptr<float[]>& in_ptr2, std::unique_ptr<int[]>& in_ptr2_dims, int in_ptr2_dims_size, int ptr2_size,  int kernel_cols, int kernel_rows, int stride_cols, int stride_rows) {
    int ptr1_cols = in_ptr1_dims[0]; // This is the full size unkerneled
    int ptr1_rows = in_ptr1_dims[1];
    int ptr2_cols = in_ptr2_dims[0]; // This is the kernel size!
    int ptr2_rows = in_ptr2_dims[1];
    int depths = 1;
    for (int i = 2; i < in_ptr1_dims_size; i++) {
        depths *= in_ptr1_dims[i];
    }

    int gpu_ptr1_bytes = ptr1_size * sizeof(float);
    int gpu_ptr2_bytes = ptr2_size * sizeof(float);

    float* gpu_ptr1;
    float* gpu_ptr2;
    float* gpu_ptr3;
    cudaMalloc(&gpu_ptr1, gpu_ptr1_bytes);
    cudaMalloc(&gpu_ptr2, gpu_ptr2_bytes);
    cudaMalloc(&gpu_ptr3, gpu_ptr1_bytes);
    cudaMemcpy(gpu_ptr1, in_ptr1.get(), gpu_ptr1_bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(gpu_ptr2, in_ptr2.get(), gpu_ptr2_bytes, cudaMemcpyHostToDevice);

    // Now what memory blocks are we going to use for this?
    int grid_cols = (ptr1_cols + std::sqrt(THREAD_SIZE_XY / THREAD_SIZE_Z) - 1) / std::sqrt(THREAD_SIZE_XY / THREAD_SIZE_Z);
    int grid_rows = (ptr1_rows + std::sqrt(THREAD_SIZE_XY / THREAD_SIZE_Z) - 1) / std::sqrt(THREAD_SIZE_XY / THREAD_SIZE_Z);
    int grid_depths = (depths + THREAD_SIZE_Z - 1) / THREAD_SIZE_Z;

    dim3 gridSize(grid_cols, grid_cols, grid_depths);
    dim3 threadSize(std::sqrt(THREAD_SIZE_XY / THREAD_SIZE_Z), std::sqrt(THREAD_SIZE_XY / THREAD_SIZE_Z), THREAD_SIZE_Z);

    poolingDerivD <<< gridSize, threadSize >>> (ptr1_cols, ptr1_rows, depths, kernel_cols, kernel_rows, stride_cols, stride_rows, gpu_ptr1, gpu_ptr2, gpu_ptr3);

    std::unique_ptr<float[]> out_ptr(new float[ptr1_size]);
    cudaMemcpy(out_ptr.get(), gpu_ptr3, gpu_ptr1_bytes, cudaMemcpyDeviceToHost);

    cudaFree(gpu_ptr1);
    cudaFree(gpu_ptr2);
    cudaFree(gpu_ptr3);

    return out_ptr;
}

__global__
void dupeD(int cols, int rows, int depths, int duped_depths, float* ptr1, float* ptr2) {
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int depth = blockIdx.z * blockDim.z + threadIdx.z; // Now represents the depth of the unstreteched size

    if ((col < cols) && (row < rows) && (depth < duped_depths)) {
        int ptr1_depth = depth % depths;
        ptr2[depth * rows * cols + row * cols + col] = ptr1[ptr1_depth * rows * cols + row * cols + col];
    }
}

// This is the broken function
std::unique_ptr<float[]> CUDAdupe(std::unique_ptr<float[]>& in_ptr1, std::unique_ptr<int[]>& in_ptr1_dims, int in_ptr1_dims_size, int ptr1_size, int dupe_size) {
    int ptr1_cols = in_ptr1_dims[0];
    int ptr1_rows = in_ptr1_dims[1];
    int depths = 1;
    for (int i = 2; i < in_ptr1_dims_size; i++) {
        depths *= in_ptr1_dims[i];
    }

    int ptr2_depths = dupe_size * depths;
    int ptr2_size = ptr1_cols * ptr1_rows * ptr2_depths;

    int gpu_ptr1_bytes = ptr1_size * sizeof(float);
    int gpu_ptr2_bytes = ptr2_size * sizeof(float);

    float* gpu_ptr1;
    float* gpu_ptr2;
    cudaMalloc(&gpu_ptr1, gpu_ptr1_bytes);
    cudaMalloc(&gpu_ptr2, gpu_ptr2_bytes);
    cudaMemcpy(gpu_ptr1, in_ptr1.get(), gpu_ptr1_bytes, cudaMemcpyHostToDevice);

    int grid_cols = (ptr1_cols + std::sqrt(THREAD_SIZE_XY / THREAD_SIZE_Z) - 1) / std::sqrt(THREAD_SIZE_XY / THREAD_SIZE_Z);
    int grid_rows = (ptr1_rows + std::sqrt(THREAD_SIZE_XY / THREAD_SIZE_Z) - 1) / std::sqrt(THREAD_SIZE_XY / THREAD_SIZE_Z);
    int grid_depths = (ptr2_depths + THREAD_SIZE_Z - 1) / THREAD_SIZE_Z; // This should be the depths of the duped one

    dim3 gridSize(grid_cols, grid_cols, grid_depths);
    dim3 threadSize(std::sqrt(THREAD_SIZE_XY / THREAD_SIZE_Z), std::sqrt(THREAD_SIZE_XY / THREAD_SIZE_Z), THREAD_SIZE_Z);

    dupeD <<< gridSize, threadSize >>> (ptr1_cols, ptr1_rows, depths, ptr2_depths, gpu_ptr1, gpu_ptr2);

    std::unique_ptr<float[]> out_ptr(new float[ptr2_size]);
    cudaMemcpy(out_ptr.get(), gpu_ptr2, gpu_ptr2_bytes, cudaMemcpyDeviceToHost);

    cudaFree(gpu_ptr1);
    cudaFree(gpu_ptr2);

    return out_ptr;
}

__global__
// What row, col and depth are we choosing? The big one
void convolutionD(int cols, int rows, int kernel_cols, int kernel_rows, int depths, int stride_cols, int stride_rows, float* ptr1, float* ptr2, float* ptr3) {
    // In here we take the correct stride and perform the convolution over that desired block for each element
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int depth = blockIdx.z * blockDim.z + threadIdx.z;

    if ((col < cols - kernel_cols + 1) && (row < rows - kernel_rows + 1) && (depth < depths)) {

        if ((col % stride_cols == 0) && (row % stride_rows == 0)) {

            float weighted = 0;

            for (int i = 0; i < kernel_rows; i++) {
                for (int j = 0; j < kernel_cols; j++) {

                    weighted += ptr1[depth * rows * cols + (row + i) * cols + (col + j)] * ptr2[depth * rows * cols + i * kernel_cols + j]; // Now I have to do the dot product of the kernel and the convolved

                }
            }

            int weighted_cols_size = (cols - kernel_cols + stride_cols) / stride_cols;
            int weighted_rows_size = (rows - kernel_rows + stride_rows) / stride_rows;

            int weighted_col = (col - kernel_cols + stride_cols) / stride_cols;
            if (weighted_col < 0) weighted_col = 0;
            int weighted_row = (row - kernel_rows + stride_rows) / stride_rows;
            if (weighted_row < 0) weighted_row = 0;            

            ptr3[depth * weighted_rows_size * weighted_cols_size + weighted_row * weighted_cols_size + weighted_col] = weighted;

        }

    }
}

// No bias is required for this
// std::unique_ptr<float[]> CUDAconvolution(std::unique_ptr<float[]>& in_ptr1, std::unique_ptr<int[]>& in_ptr1_dims, int in_ptr1_dims_size, int ptr1_size, std::unique_ptr<float[]>& in_ptr2, std::unique_ptr<int[]>& in_ptr2_dims, int in_ptr2_dims_size, int ptr2_size, int stride_cols, int stride_rows) {

//     // Convolve layer
//     int ptr1_cols = in_ptr1_dims[0];
//     int ptr1_rows = in_ptr1_dims[1];
//     int ptr1_depths = 1;
//     for (int i = 0; i < in_ptr1_dims_size; i++) { 
//         ptr1_depths *= in_ptr1_dims[i];
//     }

//     // Kernel
//     int ptr2_cols = in_ptr2_dims[0];
//     int ptr2_rows = in_ptr2_dims[1];
//     int ptr2_depths = 1;
//     for (int i = 0; i < in_ptr2_dims_size; i++) { 
//         ptr2_depths *= in_ptr2_dims[i];
//     }

//     // This will be the amount to scale the pointers for its depth size
//     int dupe_ptr1 = 1; 
//     if (in_ptr2_dims_size > 3) dupe_ptr1 = in_ptr2_dims[3];
//     int dupe_ptr2 = 1;
//     if (in_ptr1_dims_size > 3) dupe_ptr2 = in_ptr1_dims[3];

//     // We see that the dupe function duplicates every depth in each fourth dimension
//     std::unique_ptr<float[]> ptr1_duped = CUDAdupe(in_ptr1, in_ptr1_dims, in_ptr1_dims_size, ptr1_size, dupe_ptr1); // This will be the ptr1 that has been scaled to match the filter sizes
//     std::unique_ptr<float[]> ptr2_duped = CUDAdupe(in_ptr2, in_ptr2_dims, in_ptr2_dims_size, ptr2_size, dupe_ptr2); // This will scale the kernel to match the amount of input blocks there are

//     int ptr1_duped_size = ptr1_size * dupe_ptr1;
//     int ptr2_duped_size = ptr2_size * dupe_ptr2; // This part could be the problem?

//     // This part is all safe
//     int ptr3_cols = (ptr1_cols - ptr2_cols + stride_cols) / stride_cols;
//     int ptr3_rows = (ptr1_rows - ptr2_rows + stride_rows) / stride_rows;
//     int ptr3_depths = dupe_ptr1 * ptr1_depths;
//     int ptr3_size = ptr3_depths * ptr3_rows * ptr3_cols;

//     int gpu_ptr1_bytes = ptr1_duped_size * sizeof(float); // These must be the wrong allocation sizes
//     int gpu_ptr2_bytes = ptr2_duped_size * sizeof(float);
//     int gpu_ptr3_bytes = ptr3_size * sizeof(float);

//     float* gpu_ptr1; // Convolved
//     float* gpu_ptr2; // Kernel
//     float* gpu_ptr3; // Output
//     cudaMalloc(&gpu_ptr1, gpu_ptr1_bytes);
//     cudaMalloc(&gpu_ptr2, gpu_ptr2_bytes);
//     cudaMalloc(&gpu_ptr3, gpu_ptr3_bytes);
//     cudaMemcpy(gpu_ptr1, ptr1_duped.get(), gpu_ptr1_bytes, cudaMemcpyHostToDevice);
//     cudaMemcpy(gpu_ptr2, ptr2_duped.get(), gpu_ptr2_bytes, cudaMemcpyHostToDevice); // The memory allocation for this one is wrong

//     int grid_cols = (ptr3_cols + std::sqrt(THREAD_SIZE_XY / THREAD_SIZE_Z) - 1) / std::sqrt(THREAD_SIZE_XY / THREAD_SIZE_Z);
//     int grid_rows = (ptr3_rows + std::sqrt(THREAD_SIZE_XY / THREAD_SIZE_Z) - 1) / std::sqrt(THREAD_SIZE_XY / THREAD_SIZE_Z);
//     int grid_depths = (ptr3_depths + THREAD_SIZE_Z - 1) / THREAD_SIZE_Z;

//     dim3 gridSize(grid_cols, grid_cols, grid_depths);
//     dim3 threadSize(std::sqrt(THREAD_SIZE_XY / THREAD_SIZE_Z), std::sqrt(THREAD_SIZE_XY / THREAD_SIZE_Z), THREAD_SIZE_Z);

//     std::unique_ptr<float[]> out_ptr(new float[ptr3_size]);
//     cudaMemcpy(out_ptr.get(), gpu_ptr3, gpu_ptr3_bytes, cudaMemcpyDeviceToHost);

//     cudaFree(gpu_ptr1);
//     cudaFree(gpu_ptr2);
//     cudaFree(gpu_ptr3);

//     return out_ptr;
// }

// std::unique_ptr<float[]> CUDAconvolution(std::unique_ptr<float[]>& in_ptr1, std::unique_ptr<int[]>& in_ptr1_dims, int in_ptr1_dims_size, int ptr1_size, std::unique_ptr<float[]>& in_ptr2, std::unique_ptr<int[]>& in_ptr2_dims, int in_ptr2_dims_size, int ptr2_size, int stride_cols, int stride_rows) {
//     int ptr1_cols = ;
// }

// New Pseudo:
//  Inputs: A layered 4 dimensional input block
//          A layered 4 dimensional weight block (with the same depth as those of the filters)
//          The third dimensions nof each should line up but not the fourth dimension

//  Scaling: For the input block scale them to be the same size as the fourth dimension of the weight block 
//           For the weight blocks, condense it all into a single 3d layer and then scale them by the fourth dimension of the input block

// Post Scaling: Turn the scaled input block into a single three dimensional layer (do this by multiplying the depth by the rest of the size)
//               Turn the scaled weight block into a big single three dimensional block too

// Remaining steps: Perform the convolution across every different subsection
// Output it as a block with dimensions of the new rows and cols, the depth of the original depth of the input block and the fourth dimension of the kernels 

// Post processing ----------- (NOT NEEDED)
//  Do the sum across all of the fourth dimensions into a single third dimension (or something)
//  Add the bias term to each respective element

// Thoughts?
// How would this deal with a block size larger than four dimensions?
// To do so it appears that the dupe function is broken - It does not perform the duplicates properly for just the fourth dimensions, lets check this out
	.file	"tmpxft_002a92e8_00000000-6_tensoroperations.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB4316:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE4316:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.type	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0, @function
_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0:
.LFB5076:
	.cfi_startproc
	jmp	cudaLaunchKernel@PLT
	.cfi_endproc
.LFE5076:
	.size	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0, .-_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	.type	_Z10cudaMallocIfE9cudaErrorPPT_m.isra.0, @function
_Z10cudaMallocIfE9cudaErrorPPT_m.isra.0:
.LFB5077:
	.cfi_startproc
	jmp	cudaMalloc@PLT
	.cfi_endproc
.LFE5077:
	.size	_Z10cudaMallocIfE9cudaErrorPPT_m.isra.0, .-_Z10cudaMallocIfE9cudaErrorPPT_m.isra.0
	.globl	_Z40__device_stub__Z15addElementwiseDiPfS_S_iPfS_S_
	.type	_Z40__device_stub__Z15addElementwiseDiPfS_S_iPfS_S_, @function
_Z40__device_stub__Z15addElementwiseDiPfS_S_iPfS_S_:
.LFB4338:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movl	%edi, 28(%rsp)
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	32(%rsp), %rdx
	movq	%rcx, (%rsp)
	leaq	40(%rsp), %rcx
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	28(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 72(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 80(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 88(%rsp)
	movq	%rsp, %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L4
	pushq	40(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z15addElementwiseDiPfS_S_(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 144
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	88(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L4:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L6
	call	__stack_chk_fail@PLT
.L6:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4338:
	.size	_Z40__device_stub__Z15addElementwiseDiPfS_S_iPfS_S_, .-_Z40__device_stub__Z15addElementwiseDiPfS_S_iPfS_S_
	.globl	_Z15addElementwiseDiPfS_S_
	.type	_Z15addElementwiseDiPfS_S_, @function
_Z15addElementwiseDiPfS_S_:
.LFB4339:
	.cfi_startproc
	endbr64
	jmp	_Z40__device_stub__Z15addElementwiseDiPfS_S_iPfS_S_
	.cfi_endproc
.LFE4339:
	.size	_Z15addElementwiseDiPfS_S_, .-_Z15addElementwiseDiPfS_S_
	.globl	_Z45__device_stub__Z20subtractElementwiseDiPfS_S_iPfS_S_
	.type	_Z45__device_stub__Z20subtractElementwiseDiPfS_S_iPfS_S_, @function
_Z45__device_stub__Z20subtractElementwiseDiPfS_S_iPfS_S_:
.LFB4340:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movl	%edi, 28(%rsp)
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	32(%rsp), %rdx
	movq	%rcx, (%rsp)
	leaq	40(%rsp), %rcx
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	28(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 72(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 80(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 88(%rsp)
	movq	%rsp, %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L10
	pushq	40(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z20subtractElementwiseDiPfS_S_(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 144
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	88(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L10:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L12
	call	__stack_chk_fail@PLT
.L12:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4340:
	.size	_Z45__device_stub__Z20subtractElementwiseDiPfS_S_iPfS_S_, .-_Z45__device_stub__Z20subtractElementwiseDiPfS_S_iPfS_S_
	.globl	_Z20subtractElementwiseDiPfS_S_
	.type	_Z20subtractElementwiseDiPfS_S_, @function
_Z20subtractElementwiseDiPfS_S_:
.LFB4341:
	.cfi_startproc
	endbr64
	jmp	_Z45__device_stub__Z20subtractElementwiseDiPfS_S_iPfS_S_
	.cfi_endproc
.LFE4341:
	.size	_Z20subtractElementwiseDiPfS_S_, .-_Z20subtractElementwiseDiPfS_S_
	.globl	_Z45__device_stub__Z20multiplyElementwiseDiPfS_S_iPfS_S_
	.type	_Z45__device_stub__Z20multiplyElementwiseDiPfS_S_iPfS_S_, @function
_Z45__device_stub__Z20multiplyElementwiseDiPfS_S_iPfS_S_:
.LFB4342:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movl	%edi, 28(%rsp)
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	32(%rsp), %rdx
	movq	%rcx, (%rsp)
	leaq	40(%rsp), %rcx
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	28(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 72(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 80(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 88(%rsp)
	movq	%rsp, %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L15
	pushq	40(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z20multiplyElementwiseDiPfS_S_(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 144
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	88(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L15:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L17
	call	__stack_chk_fail@PLT
.L17:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4342:
	.size	_Z45__device_stub__Z20multiplyElementwiseDiPfS_S_iPfS_S_, .-_Z45__device_stub__Z20multiplyElementwiseDiPfS_S_iPfS_S_
	.globl	_Z20multiplyElementwiseDiPfS_S_
	.type	_Z20multiplyElementwiseDiPfS_S_, @function
_Z20multiplyElementwiseDiPfS_S_:
.LFB4343:
	.cfi_startproc
	endbr64
	jmp	_Z45__device_stub__Z20multiplyElementwiseDiPfS_S_iPfS_S_
	.cfi_endproc
.LFE4343:
	.size	_Z20multiplyElementwiseDiPfS_S_, .-_Z20multiplyElementwiseDiPfS_S_
	.globl	_Z43__device_stub__Z18divideElementwiseDiPfS_S_iPfS_S_
	.type	_Z43__device_stub__Z18divideElementwiseDiPfS_S_iPfS_S_, @function
_Z43__device_stub__Z18divideElementwiseDiPfS_S_iPfS_S_:
.LFB4344:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movl	%edi, 28(%rsp)
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	32(%rsp), %rdx
	movq	%rcx, (%rsp)
	leaq	40(%rsp), %rcx
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	28(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 72(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 80(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 88(%rsp)
	movq	%rsp, %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L20
	pushq	40(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z18divideElementwiseDiPfS_S_(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 144
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	88(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L20:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L22
	call	__stack_chk_fail@PLT
.L22:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4344:
	.size	_Z43__device_stub__Z18divideElementwiseDiPfS_S_iPfS_S_, .-_Z43__device_stub__Z18divideElementwiseDiPfS_S_iPfS_S_
	.globl	_Z18divideElementwiseDiPfS_S_
	.type	_Z18divideElementwiseDiPfS_S_, @function
_Z18divideElementwiseDiPfS_S_:
.LFB4345:
	.cfi_startproc
	endbr64
	jmp	_Z43__device_stub__Z18divideElementwiseDiPfS_S_iPfS_S_
	.cfi_endproc
.LFE4345:
	.size	_Z18divideElementwiseDiPfS_S_, .-_Z18divideElementwiseDiPfS_S_
	.globl	_Z42__device_stub__Z17powerElementwiseDiPfS_S_iPfS_S_
	.type	_Z42__device_stub__Z17powerElementwiseDiPfS_S_iPfS_S_, @function
_Z42__device_stub__Z17powerElementwiseDiPfS_S_iPfS_S_:
.LFB4346:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movl	%edi, 28(%rsp)
	leaq	48(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	60(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	32(%rsp), %rdx
	movq	%rcx, (%rsp)
	leaq	40(%rsp), %rcx
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	28(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 72(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 80(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 88(%rsp)
	movq	%rsp, %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L25
	pushq	40(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z17powerElementwiseDiPfS_S_(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 144
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	88(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L25:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L27
	call	__stack_chk_fail@PLT
.L27:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4346:
	.size	_Z42__device_stub__Z17powerElementwiseDiPfS_S_iPfS_S_, .-_Z42__device_stub__Z17powerElementwiseDiPfS_S_iPfS_S_
	.globl	_Z17powerElementwiseDiPfS_S_
	.type	_Z17powerElementwiseDiPfS_S_, @function
_Z17powerElementwiseDiPfS_S_:
.LFB4347:
	.cfi_startproc
	endbr64
	jmp	_Z42__device_stub__Z17powerElementwiseDiPfS_S_iPfS_S_
	.cfi_endproc
.LFE4347:
	.size	_Z17powerElementwiseDiPfS_S_, .-_Z17powerElementwiseDiPfS_S_
	.globl	_Z35__device_stub__Z10transposeDiiiPfS_iiiPfS_
	.type	_Z35__device_stub__Z10transposeDiiiPfS_iiiPfS_, @function
_Z35__device_stub__Z10transposeDiiiPfS_iiiPfS_:
.LFB4348:
	.cfi_startproc
	endbr64
	subq	$136, %rsp
	.cfi_def_cfa_offset 144
	movl	%edi, 28(%rsp)
	leaq	56(%rsp), %rdi
	movl	%esi, 24(%rsp)
	leaq	68(%rsp), %rsi
	movl	%edx, 20(%rsp)
	leaq	40(%rsp), %rdx
	movq	%rcx, 8(%rsp)
	leaq	48(%rsp), %rcx
	movq	%r8, (%rsp)
	movq	%fs:40, %rax
	movq	%rax, 120(%rsp)
	xorl	%eax, %eax
	leaq	28(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 80(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 88(%rsp)
	leaq	20(%rsp), %rax
	movq	%rax, 96(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 104(%rsp)
	movq	%rsp, %rax
	movq	%rax, 112(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 56(%rsp)
	movq	%rax, 68(%rsp)
	movl	$1, 76(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L30
	pushq	48(%rsp)
	.cfi_def_cfa_offset 152
	leaq	_Z10transposeDiiiPfS_(%rip), %rdi
	pushq	48(%rsp)
	.cfi_def_cfa_offset 160
	movq	84(%rsp), %rcx
	movl	92(%rsp), %r8d
	movq	72(%rsp), %rsi
	movl	80(%rsp), %edx
	leaq	96(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 152
	popq	%rdx
	.cfi_def_cfa_offset 144
.L30:
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L32
	call	__stack_chk_fail@PLT
.L32:
	addq	$136, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4348:
	.size	_Z35__device_stub__Z10transposeDiiiPfS_iiiPfS_, .-_Z35__device_stub__Z10transposeDiiiPfS_iiiPfS_
	.globl	_Z10transposeDiiiPfS_
	.type	_Z10transposeDiiiPfS_, @function
_Z10transposeDiiiPfS_:
.LFB4349:
	.cfi_startproc
	endbr64
	jmp	_Z35__device_stub__Z10transposeDiiiPfS_iiiPfS_
	.cfi_endproc
.LFE4349:
	.size	_Z10transposeDiiiPfS_, .-_Z10transposeDiiiPfS_
	.globl	_Z36__device_stub__Z9multiplyDiiiiPfS_S_iiiiPfS_S_
	.type	_Z36__device_stub__Z9multiplyDiiiiPfS_S_iiiiPfS_S_, @function
_Z36__device_stub__Z9multiplyDiiiiPfS_S_iiiiPfS_S_:
.LFB4350:
	.cfi_startproc
	endbr64
	subq	$168, %rsp
	.cfi_def_cfa_offset 176
	movq	176(%rsp), %rax
	movl	%edi, 44(%rsp)
	leaq	72(%rsp), %rdi
	movl	%esi, 40(%rsp)
	leaq	84(%rsp), %rsi
	movl	%edx, 36(%rsp)
	leaq	56(%rsp), %rdx
	movl	%ecx, 32(%rsp)
	leaq	64(%rsp), %rcx
	movq	%r8, 24(%rsp)
	movq	%r9, 16(%rsp)
	movq	%rax, 8(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 152(%rsp)
	xorl	%eax, %eax
	leaq	44(%rsp), %rax
	movl	$1, 80(%rsp)
	movq	%rax, 96(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	36(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 144(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 72(%rsp)
	movq	%rax, 84(%rsp)
	movl	$1, 92(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L35
	pushq	64(%rsp)
	.cfi_def_cfa_offset 184
	leaq	_Z9multiplyDiiiiPfS_S_(%rip), %rdi
	pushq	64(%rsp)
	.cfi_def_cfa_offset 192
	movq	100(%rsp), %rcx
	movl	108(%rsp), %r8d
	movq	88(%rsp), %rsi
	movl	96(%rsp), %edx
	leaq	112(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 184
	popq	%rdx
	.cfi_def_cfa_offset 176
.L35:
	movq	152(%rsp), %rax
	subq	%fs:40, %rax
	je	.L37
	call	__stack_chk_fail@PLT
.L37:
	addq	$168, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4350:
	.size	_Z36__device_stub__Z9multiplyDiiiiPfS_S_iiiiPfS_S_, .-_Z36__device_stub__Z9multiplyDiiiiPfS_S_iiiiPfS_S_
	.globl	_Z9multiplyDiiiiPfS_S_
	.type	_Z9multiplyDiiiiPfS_S_, @function
_Z9multiplyDiiiiPfS_S_:
.LFB4351:
	.cfi_startproc
	endbr64
	jmp	_Z36__device_stub__Z9multiplyDiiiiPfS_S_iiiiPfS_S_
	.cfi_endproc
.LFE4351:
	.size	_Z9multiplyDiiiiPfS_S_, .-_Z9multiplyDiiiiPfS_S_
	.globl	_Z31__device_stub__Z7rotateDiiiPfS_iiiPfS_
	.type	_Z31__device_stub__Z7rotateDiiiPfS_iiiPfS_, @function
_Z31__device_stub__Z7rotateDiiiPfS_iiiPfS_:
.LFB4352:
	.cfi_startproc
	endbr64
	subq	$136, %rsp
	.cfi_def_cfa_offset 144
	movl	%edi, 28(%rsp)
	leaq	56(%rsp), %rdi
	movl	%esi, 24(%rsp)
	leaq	68(%rsp), %rsi
	movl	%edx, 20(%rsp)
	leaq	40(%rsp), %rdx
	movq	%rcx, 8(%rsp)
	leaq	48(%rsp), %rcx
	movq	%r8, (%rsp)
	movq	%fs:40, %rax
	movq	%rax, 120(%rsp)
	xorl	%eax, %eax
	leaq	28(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 80(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 88(%rsp)
	leaq	20(%rsp), %rax
	movq	%rax, 96(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 104(%rsp)
	movq	%rsp, %rax
	movq	%rax, 112(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 56(%rsp)
	movq	%rax, 68(%rsp)
	movl	$1, 76(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L40
	pushq	48(%rsp)
	.cfi_def_cfa_offset 152
	leaq	_Z7rotateDiiiPfS_(%rip), %rdi
	pushq	48(%rsp)
	.cfi_def_cfa_offset 160
	movq	84(%rsp), %rcx
	movl	92(%rsp), %r8d
	movq	72(%rsp), %rsi
	movl	80(%rsp), %edx
	leaq	96(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 152
	popq	%rdx
	.cfi_def_cfa_offset 144
.L40:
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L42
	call	__stack_chk_fail@PLT
.L42:
	addq	$136, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4352:
	.size	_Z31__device_stub__Z7rotateDiiiPfS_iiiPfS_, .-_Z31__device_stub__Z7rotateDiiiPfS_iiiPfS_
	.globl	_Z7rotateDiiiPfS_
	.type	_Z7rotateDiiiPfS_, @function
_Z7rotateDiiiPfS_:
.LFB4353:
	.cfi_startproc
	endbr64
	jmp	_Z31__device_stub__Z7rotateDiiiPfS_iiiPfS_
	.cfi_endproc
.LFE4353:
	.size	_Z7rotateDiiiPfS_, .-_Z7rotateDiiiPfS_
	.globl	_Z40__device_stub__Z11maxPoolingDiiiiiiiPfS_iiiiiiiPfS_
	.type	_Z40__device_stub__Z11maxPoolingDiiiiiiiPfS_iiiiiiiPfS_, @function
_Z40__device_stub__Z11maxPoolingDiiiiiiiPfS_iiiiiiiPfS_:
.LFB4354:
	.cfi_startproc
	endbr64
	subq	$184, %rsp
	.cfi_def_cfa_offset 192
	movq	200(%rsp), %rax
	movl	%edi, 44(%rsp)
	leaq	72(%rsp), %rdi
	movl	%esi, 40(%rsp)
	leaq	84(%rsp), %rsi
	movq	%rax, 16(%rsp)
	movq	208(%rsp), %rax
	movl	%edx, 36(%rsp)
	leaq	56(%rsp), %rdx
	movl	%ecx, 32(%rsp)
	leaq	64(%rsp), %rcx
	movl	%r8d, 28(%rsp)
	movl	%r9d, 24(%rsp)
	movq	%rax, 8(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 168(%rsp)
	xorl	%eax, %eax
	leaq	44(%rsp), %rax
	movl	$1, 80(%rsp)
	movq	%rax, 96(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	36(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	28(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	192(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 160(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 72(%rsp)
	movq	%rax, 84(%rsp)
	movl	$1, 92(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L45
	pushq	64(%rsp)
	.cfi_def_cfa_offset 200
	leaq	_Z11maxPoolingDiiiiiiiPfS_(%rip), %rdi
	pushq	64(%rsp)
	.cfi_def_cfa_offset 208
	movq	100(%rsp), %rcx
	movl	108(%rsp), %r8d
	movq	88(%rsp), %rsi
	movl	96(%rsp), %edx
	leaq	112(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 200
	popq	%rdx
	.cfi_def_cfa_offset 192
.L45:
	movq	168(%rsp), %rax
	subq	%fs:40, %rax
	je	.L47
	call	__stack_chk_fail@PLT
.L47:
	addq	$184, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4354:
	.size	_Z40__device_stub__Z11maxPoolingDiiiiiiiPfS_iiiiiiiPfS_, .-_Z40__device_stub__Z11maxPoolingDiiiiiiiPfS_iiiiiiiPfS_
	.globl	_Z11maxPoolingDiiiiiiiPfS_
	.type	_Z11maxPoolingDiiiiiiiPfS_, @function
_Z11maxPoolingDiiiiiiiPfS_:
.LFB4355:
	.cfi_startproc
	endbr64
	jmp	_Z40__device_stub__Z11maxPoolingDiiiiiiiPfS_iiiiiiiPfS_
	.cfi_endproc
.LFE4355:
	.size	_Z11maxPoolingDiiiiiiiPfS_, .-_Z11maxPoolingDiiiiiiiPfS_
	.globl	_Z44__device_stub__Z13poolingDerivDiiiiiiiPfS_S_iiiiiiiPfS_S_
	.type	_Z44__device_stub__Z13poolingDerivDiiiiiiiPfS_S_iiiiiiiPfS_S_, @function
_Z44__device_stub__Z13poolingDerivDiiiiiiiPfS_S_iiiiiiiPfS_S_:
.LFB4356:
	.cfi_startproc
	endbr64
	subq	$184, %rsp
	.cfi_def_cfa_offset 192
	movq	200(%rsp), %rax
	movl	%edi, 44(%rsp)
	leaq	64(%rsp), %rdi
	movl	%esi, 40(%rsp)
	leaq	76(%rsp), %rsi
	movq	%rax, 16(%rsp)
	movq	208(%rsp), %rax
	movl	%edx, 36(%rsp)
	leaq	48(%rsp), %rdx
	movq	%rax, 8(%rsp)
	movq	216(%rsp), %rax
	movl	%ecx, 32(%rsp)
	leaq	56(%rsp), %rcx
	movl	%r8d, 28(%rsp)
	movl	%r9d, 24(%rsp)
	movq	%rax, (%rsp)
	movq	%fs:40, %rax
	movq	%rax, 168(%rsp)
	xorl	%eax, %eax
	leaq	44(%rsp), %rax
	movl	$1, 72(%rsp)
	movq	%rax, 88(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 96(%rsp)
	leaq	36(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	28(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	192(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 152(%rsp)
	movq	%rsp, %rax
	movq	%rax, 160(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 64(%rsp)
	movq	%rax, 76(%rsp)
	movl	$1, 84(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L50
	pushq	56(%rsp)
	.cfi_def_cfa_offset 200
	leaq	_Z13poolingDerivDiiiiiiiPfS_S_(%rip), %rdi
	pushq	56(%rsp)
	.cfi_def_cfa_offset 208
	movq	92(%rsp), %rcx
	movl	100(%rsp), %r8d
	movq	80(%rsp), %rsi
	movl	88(%rsp), %edx
	leaq	104(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 200
	popq	%rdx
	.cfi_def_cfa_offset 192
.L50:
	movq	168(%rsp), %rax
	subq	%fs:40, %rax
	je	.L52
	call	__stack_chk_fail@PLT
.L52:
	addq	$184, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4356:
	.size	_Z44__device_stub__Z13poolingDerivDiiiiiiiPfS_S_iiiiiiiPfS_S_, .-_Z44__device_stub__Z13poolingDerivDiiiiiiiPfS_S_iiiiiiiPfS_S_
	.globl	_Z13poolingDerivDiiiiiiiPfS_S_
	.type	_Z13poolingDerivDiiiiiiiPfS_S_, @function
_Z13poolingDerivDiiiiiiiPfS_S_:
.LFB4357:
	.cfi_startproc
	endbr64
	jmp	_Z44__device_stub__Z13poolingDerivDiiiiiiiPfS_S_iiiiiiiPfS_S_
	.cfi_endproc
.LFE4357:
	.size	_Z13poolingDerivDiiiiiiiPfS_S_, .-_Z13poolingDerivDiiiiiiiPfS_S_
	.globl	_Z30__device_stub__Z5dupeDiiiiPfS_iiiiPfS_
	.type	_Z30__device_stub__Z5dupeDiiiiPfS_iiiiPfS_, @function
_Z30__device_stub__Z5dupeDiiiiPfS_iiiiPfS_:
.LFB4358:
	.cfi_startproc
	endbr64
	subq	$136, %rsp
	.cfi_def_cfa_offset 144
	movl	%edi, 28(%rsp)
	leaq	48(%rsp), %rdi
	movl	%esi, 24(%rsp)
	leaq	60(%rsp), %rsi
	movl	%edx, 20(%rsp)
	leaq	32(%rsp), %rdx
	movl	%ecx, 16(%rsp)
	leaq	40(%rsp), %rcx
	movq	%r8, 8(%rsp)
	movq	%r9, (%rsp)
	movq	%fs:40, %rax
	movq	%rax, 120(%rsp)
	xorl	%eax, %eax
	leaq	28(%rsp), %rax
	movl	$1, 56(%rsp)
	movq	%rax, 72(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 80(%rsp)
	leaq	20(%rsp), %rax
	movq	%rax, 88(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 96(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 104(%rsp)
	movq	%rsp, %rax
	movq	%rax, 112(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 48(%rsp)
	movq	%rax, 60(%rsp)
	movl	$1, 68(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L55
	pushq	40(%rsp)
	.cfi_def_cfa_offset 152
	leaq	_Z5dupeDiiiiPfS_(%rip), %rdi
	pushq	40(%rsp)
	.cfi_def_cfa_offset 160
	movq	76(%rsp), %rcx
	movl	84(%rsp), %r8d
	movq	64(%rsp), %rsi
	movl	72(%rsp), %edx
	leaq	88(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 152
	popq	%rdx
	.cfi_def_cfa_offset 144
.L55:
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L57
	call	__stack_chk_fail@PLT
.L57:
	addq	$136, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4358:
	.size	_Z30__device_stub__Z5dupeDiiiiPfS_iiiiPfS_, .-_Z30__device_stub__Z5dupeDiiiiPfS_iiiiPfS_
	.globl	_Z5dupeDiiiiPfS_
	.type	_Z5dupeDiiiiPfS_, @function
_Z5dupeDiiiiPfS_:
.LFB4359:
	.cfi_startproc
	endbr64
	jmp	_Z30__device_stub__Z5dupeDiiiiPfS_iiiiPfS_
	.cfi_endproc
.LFE4359:
	.size	_Z5dupeDiiiiPfS_, .-_Z5dupeDiiiiPfS_
	.globl	_Z43__device_stub__Z12convolutionDiiiiiiiPfS_S_iiiiiiiPfS_S_
	.type	_Z43__device_stub__Z12convolutionDiiiiiiiPfS_S_iiiiiiiPfS_S_, @function
_Z43__device_stub__Z12convolutionDiiiiiiiPfS_S_iiiiiiiPfS_S_:
.LFB4360:
	.cfi_startproc
	endbr64
	subq	$184, %rsp
	.cfi_def_cfa_offset 192
	movq	200(%rsp), %rax
	movl	%edi, 44(%rsp)
	leaq	64(%rsp), %rdi
	movl	%esi, 40(%rsp)
	leaq	76(%rsp), %rsi
	movq	%rax, 16(%rsp)
	movq	208(%rsp), %rax
	movl	%edx, 36(%rsp)
	leaq	48(%rsp), %rdx
	movq	%rax, 8(%rsp)
	movq	216(%rsp), %rax
	movl	%ecx, 32(%rsp)
	leaq	56(%rsp), %rcx
	movl	%r8d, 28(%rsp)
	movl	%r9d, 24(%rsp)
	movq	%rax, (%rsp)
	movq	%fs:40, %rax
	movq	%rax, 168(%rsp)
	xorl	%eax, %eax
	leaq	44(%rsp), %rax
	movl	$1, 72(%rsp)
	movq	%rax, 88(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 96(%rsp)
	leaq	36(%rsp), %rax
	movq	%rax, 104(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 112(%rsp)
	leaq	28(%rsp), %rax
	movq	%rax, 120(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 128(%rsp)
	leaq	192(%rsp), %rax
	movq	%rax, 136(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 152(%rsp)
	movq	%rsp, %rax
	movq	%rax, 160(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 64(%rsp)
	movq	%rax, 76(%rsp)
	movl	$1, 84(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L60
	pushq	56(%rsp)
	.cfi_def_cfa_offset 200
	leaq	_Z12convolutionDiiiiiiiPfS_S_(%rip), %rdi
	pushq	56(%rsp)
	.cfi_def_cfa_offset 208
	movq	92(%rsp), %rcx
	movl	100(%rsp), %r8d
	movq	80(%rsp), %rsi
	movl	88(%rsp), %edx
	leaq	104(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 200
	popq	%rdx
	.cfi_def_cfa_offset 192
.L60:
	movq	168(%rsp), %rax
	subq	%fs:40, %rax
	je	.L62
	call	__stack_chk_fail@PLT
.L62:
	addq	$184, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4360:
	.size	_Z43__device_stub__Z12convolutionDiiiiiiiPfS_S_iiiiiiiPfS_S_, .-_Z43__device_stub__Z12convolutionDiiiiiiiPfS_S_iiiiiiiPfS_S_
	.globl	_Z12convolutionDiiiiiiiPfS_S_
	.type	_Z12convolutionDiiiiiiiPfS_S_, @function
_Z12convolutionDiiiiiiiPfS_S_:
.LFB4361:
	.cfi_startproc
	endbr64
	jmp	_Z43__device_stub__Z12convolutionDiiiiiiiPfS_S_iiiiiiiPfS_S_
	.cfi_endproc
.LFE4361:
	.size	_Z12convolutionDiiiiiiiPfS_S_, .-_Z12convolutionDiiiiiiiPfS_S_
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"_Z12convolutionDiiiiiiiPfS_S_"
.LC1:
	.string	"_Z5dupeDiiiiPfS_"
.LC2:
	.string	"_Z13poolingDerivDiiiiiiiPfS_S_"
.LC3:
	.string	"_Z11maxPoolingDiiiiiiiPfS_"
.LC4:
	.string	"_Z7rotateDiiiPfS_"
.LC5:
	.string	"_Z9multiplyDiiiiPfS_S_"
.LC6:
	.string	"_Z10transposeDiiiPfS_"
.LC7:
	.string	"_Z17powerElementwiseDiPfS_S_"
.LC8:
	.string	"_Z18divideElementwiseDiPfS_S_"
.LC9:
	.string	"_Z20multiplyElementwiseDiPfS_S_"
.LC10:
	.string	"_Z20subtractElementwiseDiPfS_S_"
.LC11:
	.string	"_Z15addElementwiseDiPfS_S_"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB4363:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	pushq	$0
	.cfi_def_cfa_offset 24
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC0(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z12convolutionDiiiiiiiPfS_S_(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC1(%rip), %rdx
	movq	%rbx, %rdi
	leaq	_Z5dupeDiiiiPfS_(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC2(%rip), %rdx
	movq	%rbx, %rdi
	leaq	_Z13poolingDerivDiiiiiiiPfS_S_(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC3(%rip), %rdx
	movq	%rbx, %rdi
	leaq	_Z11maxPoolingDiiiiiiiPfS_(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC4(%rip), %rdx
	movq	%rbx, %rdi
	leaq	_Z7rotateDiiiPfS_(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC5(%rip), %rdx
	movq	%rbx, %rdi
	leaq	_Z9multiplyDiiiiPfS_S_(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC6(%rip), %rdx
	movq	%rbx, %rdi
	leaq	_Z10transposeDiiiPfS_(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC7(%rip), %rdx
	movq	%rbx, %rdi
	leaq	_Z17powerElementwiseDiPfS_S_(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	leaq	.LC8(%rip), %rdx
	pushq	$0
	.cfi_def_cfa_offset 24
	movq	%rdx, %rcx
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	_Z18divideElementwiseDiPfS_S_(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC9(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z20multiplyElementwiseDiPfS_S_(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC10(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z20subtractElementwiseDiPfS_S_(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	xorl	%r9d, %r9d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC11(%rip), %rdx
	orl	$-1, %r8d
	leaq	_Z15addElementwiseDiPfS_S_(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE4363:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.text._ZNSt10unique_ptrIA_fSt14default_deleteIS0_EED2Ev,"axG",@progbits,_ZNSt10unique_ptrIA_fSt14default_deleteIS0_EED5Ev,comdat
	.align 2
	.weak	_ZNSt10unique_ptrIA_fSt14default_deleteIS0_EED2Ev
	.type	_ZNSt10unique_ptrIA_fSt14default_deleteIS0_EED2Ev, @function
_ZNSt10unique_ptrIA_fSt14default_deleteIS0_EED2Ev:
.LFB4680:
	.cfi_startproc
	endbr64
	movq	(%rdi), %rdi
	testq	%rdi, %rdi
	je	.L67
	jmp	_ZdaPv@PLT
.L67:
	ret
	.cfi_endproc
.LFE4680:
	.size	_ZNSt10unique_ptrIA_fSt14default_deleteIS0_EED2Ev, .-_ZNSt10unique_ptrIA_fSt14default_deleteIS0_EED2Ev
	.weak	_ZNSt10unique_ptrIA_fSt14default_deleteIS0_EED1Ev
	.set	_ZNSt10unique_ptrIA_fSt14default_deleteIS0_EED1Ev,_ZNSt10unique_ptrIA_fSt14default_deleteIS0_EED2Ev
	.text
	.globl	_Z18CUDAaddElementwiseRSt10unique_ptrIA_fSt14default_deleteIS0_EES4_i
	.type	_Z18CUDAaddElementwiseRSt10unique_ptrIA_fSt14default_deleteIS0_EES4_i, @function
_Z18CUDAaddElementwiseRSt10unique_ptrIA_fSt14default_deleteIS0_EES4_i:
.LFB4288:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA4288
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movq	%rsi, %r15
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	movq	%rdx, %r14
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	movslq	%ecx, %r13
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movq	%rdi, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	leal	0(,%r13,4), %ebx
	movslq	%ebx, %rbx
	movq	%rbx, %rsi
	subq	$72, %rsp
	.cfi_def_cfa_offset 128
	movq	%fs:40, %rax
	movq	%rax, 56(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rdi
.LEHB0:
	call	_Z10cudaMallocIfE9cudaErrorPPT_m.isra.0
	leaq	16(%rsp), %rdi
	movq	%rbx, %rsi
	call	_Z10cudaMallocIfE9cudaErrorPPT_m.isra.0
	leaq	24(%rsp), %rdi
	movq	%rbx, %rsi
	call	_Z10cudaMallocIfE9cudaErrorPPT_m.isra.0
	movq	(%r15), %rsi
	movq	8(%rsp), %rdi
	movq	%rbx, %rdx
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	movq	(%r14), %rsi
	movq	16(%rsp), %rdi
	movq	%rbx, %rdx
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	movl	$1024, %ecx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	leal	1023(%r13), %eax
	cltd
	idivl	%ecx
	movl	$4194305, %edx
	movl	$1, %ecx
	salq	$10, %rdx
	movl	%eax, 32(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 36(%rsp)
	movq	32(%rsp), %rdi
	movl	40(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L70
	movq	24(%rsp), %rcx
	movq	16(%rsp), %rdx
	movl	%r13d, %edi
	movq	8(%rsp), %rsi
	call	_Z40__device_stub__Z15addElementwiseDiPfS_S_iPfS_S_
.L70:
	movabsq	$2305843009213693950, %rax
	cmpq	%r13, %rax
	jb	.L71
	leaq	0(,%r13,4), %rdi
	call	_Znam@PLT
.LEHE0:
	movq	24(%rsp), %rsi
	movl	$2, %ecx
	movq	%rbx, %rdx
	movq	%rax, (%r12)
	movq	%rax, %rdi
.LEHB1:
	call	cudaMemcpy@PLT
.LEHE1:
	jmp	.L83
.L71:
	movq	56(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L82
.LEHB2:
	call	__cxa_throw_bad_array_new_length@PLT
.LEHE2:
.L83:
	movq	8(%rsp), %rdi
.LEHB3:
	call	cudaFree@PLT
	movq	16(%rsp), %rdi
	call	cudaFree@PLT
	movq	24(%rsp), %rdi
	call	cudaFree@PLT
.LEHE3:
	jmp	.L84
.L77:
	endbr64
	movq	%rax, %rbx
.L74:
	movq	%r12, %rdi
	call	_ZNSt10unique_ptrIA_fSt14default_deleteIS0_EED1Ev
	movq	56(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L82
	movq	%rbx, %rdi
.LEHB4:
	call	_Unwind_Resume@PLT
.LEHE4:
.L84:
	movq	56(%rsp), %rax
	subq	%fs:40, %rax
	je	.L76
.L82:
	call	__stack_chk_fail@PLT
.L76:
	addq	$72, %rsp
	.cfi_def_cfa_offset 56
	movq	%r12, %rax
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4288:
	.globl	__gxx_personality_v0
	.section	.gcc_except_table,"a",@progbits
.LLSDA4288:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE4288-.LLSDACSB4288
.LLSDACSB4288:
	.uleb128 .LEHB0-.LFB4288
	.uleb128 .LEHE0-.LEHB0
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB1-.LFB4288
	.uleb128 .LEHE1-.LEHB1
	.uleb128 .L77-.LFB4288
	.uleb128 0
	.uleb128 .LEHB2-.LFB4288
	.uleb128 .LEHE2-.LEHB2
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB3-.LFB4288
	.uleb128 .LEHE3-.LEHB3
	.uleb128 .L77-.LFB4288
	.uleb128 0
	.uleb128 .LEHB4-.LFB4288
	.uleb128 .LEHE4-.LEHB4
	.uleb128 0
	.uleb128 0
.LLSDACSE4288:
	.text
	.size	_Z18CUDAaddElementwiseRSt10unique_ptrIA_fSt14default_deleteIS0_EES4_i, .-_Z18CUDAaddElementwiseRSt10unique_ptrIA_fSt14default_deleteIS0_EES4_i
	.globl	_Z23CUDAsubtractElementwiseRSt10unique_ptrIA_fSt14default_deleteIS0_EES4_i
	.type	_Z23CUDAsubtractElementwiseRSt10unique_ptrIA_fSt14default_deleteIS0_EES4_i, @function
_Z23CUDAsubtractElementwiseRSt10unique_ptrIA_fSt14default_deleteIS0_EES4_i:
.LFB4295:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA4295
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movq	%rsi, %r15
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	movq	%rdx, %r14
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	movslq	%ecx, %r13
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movq	%rdi, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	leal	0(,%r13,4), %ebx
	movslq	%ebx, %rbx
	movq	%rbx, %rsi
	subq	$72, %rsp
	.cfi_def_cfa_offset 128
	movq	%fs:40, %rax
	movq	%rax, 56(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rdi
.LEHB5:
	call	_Z10cudaMallocIfE9cudaErrorPPT_m.isra.0
	leaq	16(%rsp), %rdi
	movq	%rbx, %rsi
	call	_Z10cudaMallocIfE9cudaErrorPPT_m.isra.0
	leaq	24(%rsp), %rdi
	movq	%rbx, %rsi
	call	_Z10cudaMallocIfE9cudaErrorPPT_m.isra.0
	movq	(%r15), %rsi
	movq	8(%rsp), %rdi
	movq	%rbx, %rdx
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	movq	(%r14), %rsi
	movq	16(%rsp), %rdi
	movq	%rbx, %rdx
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	movl	$1024, %ecx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	leal	1023(%r13), %eax
	cltd
	idivl	%ecx
	movl	$4194305, %edx
	movl	$1, %ecx
	salq	$10, %rdx
	movl	%eax, 32(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 36(%rsp)
	movq	32(%rsp), %rdi
	movl	40(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L86
	movq	24(%rsp), %rcx
	movq	16(%rsp), %rdx
	movl	%r13d, %edi
	movq	8(%rsp), %rsi
	call	_Z45__device_stub__Z20subtractElementwiseDiPfS_S_iPfS_S_
.L86:
	movabsq	$2305843009213693950, %rax
	cmpq	%r13, %rax
	jb	.L87
	leaq	0(,%r13,4), %rdi
	call	_Znam@PLT
.LEHE5:
	movq	24(%rsp), %rsi
	movl	$2, %ecx
	movq	%rbx, %rdx
	movq	%rax, (%r12)
	movq	%rax, %rdi
.LEHB6:
	call	cudaMemcpy@PLT
.LEHE6:
	jmp	.L99
.L87:
	movq	56(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L98
.LEHB7:
	call	__cxa_throw_bad_array_new_length@PLT
.LEHE7:
.L99:
	movq	8(%rsp), %rdi
.LEHB8:
	call	cudaFree@PLT
	movq	16(%rsp), %rdi
	call	cudaFree@PLT
	movq	24(%rsp), %rdi
	call	cudaFree@PLT
.LEHE8:
	jmp	.L100
.L93:
	endbr64
	movq	%rax, %rbx
.L90:
	movq	%r12, %rdi
	call	_ZNSt10unique_ptrIA_fSt14default_deleteIS0_EED1Ev
	movq	56(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L98
	movq	%rbx, %rdi
.LEHB9:
	call	_Unwind_Resume@PLT
.LEHE9:
.L100:
	movq	56(%rsp), %rax
	subq	%fs:40, %rax
	je	.L92
.L98:
	call	__stack_chk_fail@PLT
.L92:
	addq	$72, %rsp
	.cfi_def_cfa_offset 56
	movq	%r12, %rax
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4295:
	.section	.gcc_except_table
.LLSDA4295:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE4295-.LLSDACSB4295
.LLSDACSB4295:
	.uleb128 .LEHB5-.LFB4295
	.uleb128 .LEHE5-.LEHB5
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB6-.LFB4295
	.uleb128 .LEHE6-.LEHB6
	.uleb128 .L93-.LFB4295
	.uleb128 0
	.uleb128 .LEHB7-.LFB4295
	.uleb128 .LEHE7-.LEHB7
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB8-.LFB4295
	.uleb128 .LEHE8-.LEHB8
	.uleb128 .L93-.LFB4295
	.uleb128 0
	.uleb128 .LEHB9-.LFB4295
	.uleb128 .LEHE9-.LEHB9
	.uleb128 0
	.uleb128 0
.LLSDACSE4295:
	.text
	.size	_Z23CUDAsubtractElementwiseRSt10unique_ptrIA_fSt14default_deleteIS0_EES4_i, .-_Z23CUDAsubtractElementwiseRSt10unique_ptrIA_fSt14default_deleteIS0_EES4_i
	.globl	_Z23CUDAmultiplyElementwiseRSt10unique_ptrIA_fSt14default_deleteIS0_EES4_i
	.type	_Z23CUDAmultiplyElementwiseRSt10unique_ptrIA_fSt14default_deleteIS0_EES4_i, @function
_Z23CUDAmultiplyElementwiseRSt10unique_ptrIA_fSt14default_deleteIS0_EES4_i:
.LFB4296:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA4296
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movq	%rsi, %r15
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	movq	%rdx, %r14
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	movslq	%ecx, %r13
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movq	%rdi, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	leal	0(,%r13,4), %ebx
	movslq	%ebx, %rbx
	movq	%rbx, %rsi
	subq	$72, %rsp
	.cfi_def_cfa_offset 128
	movq	%fs:40, %rax
	movq	%rax, 56(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rdi
.LEHB10:
	call	_Z10cudaMallocIfE9cudaErrorPPT_m.isra.0
	leaq	16(%rsp), %rdi
	movq	%rbx, %rsi
	call	_Z10cudaMallocIfE9cudaErrorPPT_m.isra.0
	leaq	24(%rsp), %rdi
	movq	%rbx, %rsi
	call	_Z10cudaMallocIfE9cudaErrorPPT_m.isra.0
	movq	(%r15), %rsi
	movq	8(%rsp), %rdi
	movq	%rbx, %rdx
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	movq	(%r14), %rsi
	movq	16(%rsp), %rdi
	movq	%rbx, %rdx
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	movl	$1024, %ecx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	leal	1023(%r13), %eax
	cltd
	idivl	%ecx
	movl	$4194305, %edx
	movl	$1, %ecx
	salq	$10, %rdx
	movl	%eax, 32(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 36(%rsp)
	movq	32(%rsp), %rdi
	movl	40(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L102
	movq	24(%rsp), %rcx
	movq	16(%rsp), %rdx
	movl	%r13d, %edi
	movq	8(%rsp), %rsi
	call	_Z45__device_stub__Z20multiplyElementwiseDiPfS_S_iPfS_S_
.L102:
	movabsq	$2305843009213693950, %rax
	cmpq	%r13, %rax
	jb	.L103
	leaq	0(,%r13,4), %rdi
	call	_Znam@PLT
.LEHE10:
	movq	24(%rsp), %rsi
	movl	$2, %ecx
	movq	%rbx, %rdx
	movq	%rax, (%r12)
	movq	%rax, %rdi
.LEHB11:
	call	cudaMemcpy@PLT
.LEHE11:
	jmp	.L115
.L103:
	movq	56(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L114
.LEHB12:
	call	__cxa_throw_bad_array_new_length@PLT
.LEHE12:
.L115:
	movq	8(%rsp), %rdi
.LEHB13:
	call	cudaFree@PLT
	movq	16(%rsp), %rdi
	call	cudaFree@PLT
	movq	24(%rsp), %rdi
	call	cudaFree@PLT
.LEHE13:
	jmp	.L116
.L109:
	endbr64
	movq	%rax, %rbx
.L106:
	movq	%r12, %rdi
	call	_ZNSt10unique_ptrIA_fSt14default_deleteIS0_EED1Ev
	movq	56(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L114
	movq	%rbx, %rdi
.LEHB14:
	call	_Unwind_Resume@PLT
.LEHE14:
.L116:
	movq	56(%rsp), %rax
	subq	%fs:40, %rax
	je	.L108
.L114:
	call	__stack_chk_fail@PLT
.L108:
	addq	$72, %rsp
	.cfi_def_cfa_offset 56
	movq	%r12, %rax
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4296:
	.section	.gcc_except_table
.LLSDA4296:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE4296-.LLSDACSB4296
.LLSDACSB4296:
	.uleb128 .LEHB10-.LFB4296
	.uleb128 .LEHE10-.LEHB10
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB11-.LFB4296
	.uleb128 .LEHE11-.LEHB11
	.uleb128 .L109-.LFB4296
	.uleb128 0
	.uleb128 .LEHB12-.LFB4296
	.uleb128 .LEHE12-.LEHB12
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB13-.LFB4296
	.uleb128 .LEHE13-.LEHB13
	.uleb128 .L109-.LFB4296
	.uleb128 0
	.uleb128 .LEHB14-.LFB4296
	.uleb128 .LEHE14-.LEHB14
	.uleb128 0
	.uleb128 0
.LLSDACSE4296:
	.text
	.size	_Z23CUDAmultiplyElementwiseRSt10unique_ptrIA_fSt14default_deleteIS0_EES4_i, .-_Z23CUDAmultiplyElementwiseRSt10unique_ptrIA_fSt14default_deleteIS0_EES4_i
	.globl	_Z21CUDAdivideElementwiseRSt10unique_ptrIA_fSt14default_deleteIS0_EES4_i
	.type	_Z21CUDAdivideElementwiseRSt10unique_ptrIA_fSt14default_deleteIS0_EES4_i, @function
_Z21CUDAdivideElementwiseRSt10unique_ptrIA_fSt14default_deleteIS0_EES4_i:
.LFB4297:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA4297
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movq	%rsi, %r15
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	movq	%rdx, %r14
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	movslq	%ecx, %r13
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movq	%rdi, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	leal	0(,%r13,4), %ebx
	movslq	%ebx, %rbx
	movq	%rbx, %rsi
	subq	$72, %rsp
	.cfi_def_cfa_offset 128
	movq	%fs:40, %rax
	movq	%rax, 56(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rdi
.LEHB15:
	call	_Z10cudaMallocIfE9cudaErrorPPT_m.isra.0
	leaq	16(%rsp), %rdi
	movq	%rbx, %rsi
	call	_Z10cudaMallocIfE9cudaErrorPPT_m.isra.0
	leaq	24(%rsp), %rdi
	movq	%rbx, %rsi
	call	_Z10cudaMallocIfE9cudaErrorPPT_m.isra.0
	movq	(%r15), %rsi
	movq	8(%rsp), %rdi
	movq	%rbx, %rdx
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	movq	(%r14), %rsi
	movq	16(%rsp), %rdi
	movq	%rbx, %rdx
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	movl	$1024, %ecx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	leal	1023(%r13), %eax
	cltd
	idivl	%ecx
	movl	$4194305, %edx
	movl	$1, %ecx
	salq	$10, %rdx
	movl	%eax, 32(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 36(%rsp)
	movq	32(%rsp), %rdi
	movl	40(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L118
	movq	24(%rsp), %rcx
	movq	16(%rsp), %rdx
	movl	%r13d, %edi
	movq	8(%rsp), %rsi
	call	_Z43__device_stub__Z18divideElementwiseDiPfS_S_iPfS_S_
.L118:
	movabsq	$2305843009213693950, %rax
	cmpq	%r13, %rax
	jb	.L119
	leaq	0(,%r13,4), %rdi
	call	_Znam@PLT
.LEHE15:
	movq	24(%rsp), %rsi
	movl	$2, %ecx
	movq	%rbx, %rdx
	movq	%rax, (%r12)
	movq	%rax, %rdi
.LEHB16:
	call	cudaMemcpy@PLT
.LEHE16:
	jmp	.L131
.L119:
	movq	56(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L130
.LEHB17:
	call	__cxa_throw_bad_array_new_length@PLT
.LEHE17:
.L131:
	movq	8(%rsp), %rdi
.LEHB18:
	call	cudaFree@PLT
	movq	16(%rsp), %rdi
	call	cudaFree@PLT
	movq	24(%rsp), %rdi
	call	cudaFree@PLT
.LEHE18:
	jmp	.L132
.L125:
	endbr64
	movq	%rax, %rbx
.L122:
	movq	%r12, %rdi
	call	_ZNSt10unique_ptrIA_fSt14default_deleteIS0_EED1Ev
	movq	56(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L130
	movq	%rbx, %rdi
.LEHB19:
	call	_Unwind_Resume@PLT
.LEHE19:
.L132:
	movq	56(%rsp), %rax
	subq	%fs:40, %rax
	je	.L124
.L130:
	call	__stack_chk_fail@PLT
.L124:
	addq	$72, %rsp
	.cfi_def_cfa_offset 56
	movq	%r12, %rax
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4297:
	.section	.gcc_except_table
.LLSDA4297:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE4297-.LLSDACSB4297
.LLSDACSB4297:
	.uleb128 .LEHB15-.LFB4297
	.uleb128 .LEHE15-.LEHB15
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB16-.LFB4297
	.uleb128 .LEHE16-.LEHB16
	.uleb128 .L125-.LFB4297
	.uleb128 0
	.uleb128 .LEHB17-.LFB4297
	.uleb128 .LEHE17-.LEHB17
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB18-.LFB4297
	.uleb128 .LEHE18-.LEHB18
	.uleb128 .L125-.LFB4297
	.uleb128 0
	.uleb128 .LEHB19-.LFB4297
	.uleb128 .LEHE19-.LEHB19
	.uleb128 0
	.uleb128 0
.LLSDACSE4297:
	.text
	.size	_Z21CUDAdivideElementwiseRSt10unique_ptrIA_fSt14default_deleteIS0_EES4_i, .-_Z21CUDAdivideElementwiseRSt10unique_ptrIA_fSt14default_deleteIS0_EES4_i
	.globl	_Z20CUDApowerElementwiseRSt10unique_ptrIA_fSt14default_deleteIS0_EES4_i
	.type	_Z20CUDApowerElementwiseRSt10unique_ptrIA_fSt14default_deleteIS0_EES4_i, @function
_Z20CUDApowerElementwiseRSt10unique_ptrIA_fSt14default_deleteIS0_EES4_i:
.LFB4298:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA4298
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movq	%rsi, %r15
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	movq	%rdx, %r14
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	movslq	%ecx, %r13
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movq	%rdi, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	leal	0(,%r13,4), %ebx
	movslq	%ebx, %rbx
	movq	%rbx, %rsi
	subq	$72, %rsp
	.cfi_def_cfa_offset 128
	movq	%fs:40, %rax
	movq	%rax, 56(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rdi
.LEHB20:
	call	_Z10cudaMallocIfE9cudaErrorPPT_m.isra.0
	leaq	16(%rsp), %rdi
	movq	%rbx, %rsi
	call	_Z10cudaMallocIfE9cudaErrorPPT_m.isra.0
	leaq	24(%rsp), %rdi
	movq	%rbx, %rsi
	call	_Z10cudaMallocIfE9cudaErrorPPT_m.isra.0
	movq	(%r15), %rsi
	movq	8(%rsp), %rdi
	movq	%rbx, %rdx
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	movq	(%r14), %rsi
	movq	16(%rsp), %rdi
	movq	%rbx, %rdx
	movl	$1, %ecx
	call	cudaMemcpy@PLT
	movl	$1024, %ecx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	leal	1023(%r13), %eax
	cltd
	idivl	%ecx
	movl	$4194305, %edx
	movl	$1, %ecx
	salq	$10, %rdx
	movl	%eax, 32(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 36(%rsp)
	movq	32(%rsp), %rdi
	movl	40(%rsp), %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L134
	movq	24(%rsp), %rcx
	movq	16(%rsp), %rdx
	movl	%r13d, %edi
	movq	8(%rsp), %rsi
	call	_Z42__device_stub__Z17powerElementwiseDiPfS_S_iPfS_S_
.L134:
	movabsq	$2305843009213693950, %rax
	cmpq	%r13, %rax
	jb	.L135
	leaq	0(,%r13,4), %rdi
	call	_Znam@PLT
.LEHE20:
	movq	24(%rsp), %rsi
	movl	$2, %ecx
	movq	%rbx, %rdx
	movq	%rax, (%r12)
	movq	%rax, %rdi
.LEHB21:
	call	cudaMemcpy@PLT
.LEHE21:
	jmp	.L147
.L135:
	movq	56(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L146
.LEHB22:
	call	__cxa_throw_bad_array_new_length@PLT
.LEHE22:
.L147:
	movq	8(%rsp), %rdi
.LEHB23:
	call	cudaFree@PLT
	movq	16(%rsp), %rdi
	call	cudaFree@PLT
	movq	24(%rsp), %rdi
	call	cudaFree@PLT
.LEHE23:
	jmp	.L148
.L141:
	endbr64
	movq	%rax, %rbx
.L138:
	movq	%r12, %rdi
	call	_ZNSt10unique_ptrIA_fSt14default_deleteIS0_EED1Ev
	movq	56(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L146
	movq	%rbx, %rdi
.LEHB24:
	call	_Unwind_Resume@PLT
.LEHE24:
.L148:
	movq	56(%rsp), %rax
	subq	%fs:40, %rax
	je	.L140
.L146:
	call	__stack_chk_fail@PLT
.L140:
	addq	$72, %rsp
	.cfi_def_cfa_offset 56
	movq	%r12, %rax
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4298:
	.section	.gcc_except_table
.LLSDA4298:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE4298-.LLSDACSB4298
.LLSDACSB4298:
	.uleb128 .LEHB20-.LFB4298
	.uleb128 .LEHE20-.LEHB20
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB21-.LFB4298
	.uleb128 .LEHE21-.LEHB21
	.uleb128 .L141-.LFB4298
	.uleb128 0
	.uleb128 .LEHB22-.LFB4298
	.uleb128 .LEHE22-.LEHB22
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB23-.LFB4298
	.uleb128 .LEHE23-.LEHB23
	.uleb128 .L141-.LFB4298
	.uleb128 0
	.uleb128 .LEHB24-.LFB4298
	.uleb128 .LEHE24-.LEHB24
	.uleb128 0
	.uleb128 0
.LLSDACSE4298:
	.text
	.size	_Z20CUDApowerElementwiseRSt10unique_ptrIA_fSt14default_deleteIS0_EES4_i, .-_Z20CUDApowerElementwiseRSt10unique_ptrIA_fSt14default_deleteIS0_EES4_i
	.globl	_Z13CUDAtransposeRSt10unique_ptrIA_fSt14default_deleteIS0_EERS_IA_iS1_IS5_EEii
	.type	_Z13CUDAtransposeRSt10unique_ptrIA_fSt14default_deleteIS0_EERS_IA_iS1_IS5_EEii, @function
_Z13CUDAtransposeRSt10unique_ptrIA_fSt14default_deleteIS0_EERS_IA_iS1_IS5_EEii:
.LFB4299:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA4299
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movl	$1, %r15d
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movq	%rdi, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movq	%rsi, %rbx
	subq	$72, %rsp
	.cfi_def_cfa_offset 128
	movq	%fs:40, %rsi
	movq	%rsi, 56(%rsp)
	movq	(%rdx), %rsi
	movl	$2, %edx
	movl	(%rsi), %r13d
	movl	4(%rsi), %r14d
.L150:
	cmpl	%edx, %ecx
	jle	.L165
	imull	(%rsi,%rdx,4), %r15d
	incq	%rdx
	jmp	.L150
.L165:
	movslq	%r8d, %r12
	sall	$2, %r8d
	leaq	16(%rsp), %rdi
	movslq	%r8d, %rax
	movq	%rax, %rsi
	movq	%rax, 8(%rsp)
.LEHB25:
	call	_Z10cudaMallocIfE9cudaErrorPPT_m.isra.0
	movq	8(%rsp), %rsi
	leaq	24(%rsp), %rdi
	call	_Z10cudaMallocIfE9cudaErrorPPT_m.isra.0
	movq	8(%rsp), %rdx
	movq	(%rbx), %rsi
	movl	$1, %ecx
	movq	16(%rsp), %rdi
	call	cudaMemcpy@PLT
	cvtsi2sdl	%r13d, %xmm0
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	addsd	.LC12(%rip), %xmm0
	subsd	.LC13(%rip), %xmm0
	movl	$64, %ecx
	mulsd	.LC14(%rip), %xmm0
	cvttsd2sil	%xmm0, %eax
	movl	%eax, 32(%rsp)
	movl	%eax, 36(%rsp)
	leal	63(%r15), %eax
	movq	32(%rsp), %rdi
	cltd
	idivl	%ecx
	movabsq	$17179869188, %rdx
	movl	$64, %ecx
	movl	%eax, %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L152
	movq	24(%rsp), %r8
	movq	16(%rsp), %rcx
	movl	%r15d, %edx
	movl	%r14d, %esi
	movl	%r13d, %edi
	call	_Z35__device_stub__Z10transposeDiiiPfS_iiiPfS_
.L152:
	movabsq	$2305843009213693950, %rax
	cmpq	%r12, %rax
	jb	.L153
	leaq	0(,%r12,4), %rdi
	call	_Znam@PLT
.LEHE25:
	movq	8(%rsp), %rdx
	movq	24(%rsp), %rsi
	movl	$2, %ecx
	movq	%rax, 0(%rbp)
	movq	%rax, %rdi
.LEHB26:
	call	cudaMemcpy@PLT
.LEHE26:
	jmp	.L166
.L153:
	movq	56(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L164
.LEHB27:
	call	__cxa_throw_bad_array_new_length@PLT
.LEHE27:
.L166:
	movq	16(%rsp), %rdi
.LEHB28:
	call	cudaFree@PLT
	movq	24(%rsp), %rdi
	call	cudaFree@PLT
.LEHE28:
	jmp	.L167
.L159:
	endbr64
	movq	%rax, %rbx
.L156:
	movq	%rbp, %rdi
	call	_ZNSt10unique_ptrIA_fSt14default_deleteIS0_EED1Ev
	movq	56(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L164
	movq	%rbx, %rdi
.LEHB29:
	call	_Unwind_Resume@PLT
.LEHE29:
.L167:
	movq	56(%rsp), %rax
	subq	%fs:40, %rax
	je	.L158
.L164:
	call	__stack_chk_fail@PLT
.L158:
	addq	$72, %rsp
	.cfi_def_cfa_offset 56
	movq	%rbp, %rax
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4299:
	.section	.gcc_except_table
.LLSDA4299:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE4299-.LLSDACSB4299
.LLSDACSB4299:
	.uleb128 .LEHB25-.LFB4299
	.uleb128 .LEHE25-.LEHB25
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB26-.LFB4299
	.uleb128 .LEHE26-.LEHB26
	.uleb128 .L159-.LFB4299
	.uleb128 0
	.uleb128 .LEHB27-.LFB4299
	.uleb128 .LEHE27-.LEHB27
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB28-.LFB4299
	.uleb128 .LEHE28-.LEHB28
	.uleb128 .L159-.LFB4299
	.uleb128 0
	.uleb128 .LEHB29-.LFB4299
	.uleb128 .LEHE29-.LEHB29
	.uleb128 0
	.uleb128 0
.LLSDACSE4299:
	.text
	.size	_Z13CUDAtransposeRSt10unique_ptrIA_fSt14default_deleteIS0_EERS_IA_iS1_IS5_EEii, .-_Z13CUDAtransposeRSt10unique_ptrIA_fSt14default_deleteIS0_EERS_IA_iS1_IS5_EEii
	.globl	_Z12CUDAmultiplyRSt10unique_ptrIA_fSt14default_deleteIS0_EERS_IA_iS1_IS5_EEiiS4_S8_ii
	.type	_Z12CUDAmultiplyRSt10unique_ptrIA_fSt14default_deleteIS0_EERS_IA_iS1_IS5_EEiiS4_S8_ii, @function
_Z12CUDAmultiplyRSt10unique_ptrIA_fSt14default_deleteIS0_EERS_IA_iS1_IS5_EEiiS4_S8_ii:
.LFB4309:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA4309
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movq	%rsi, %rax
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movl	$1, %r12d
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	subq	$120, %rsp
	.cfi_def_cfa_offset 176
	movq	%rdi, 8(%rsp)
	movq	%fs:40, %rsi
	movq	%rsi, 104(%rsp)
	movq	(%rdx), %rsi
	movq	176(%rsp), %rdx
	movq	(%rdx), %rdx
	movl	(%rdx), %r14d
	movl	$2, %edx
	movl	(%rsi), %ebx
	movl	4(%rsi), %r15d
	movl	%ebx, 20(%rsp)
.L169:
	cmpl	%edx, %ecx
	jle	.L184
	imull	(%rsi,%rdx,4), %r12d
	incq	%rdx
	jmp	.L169
.L184:
	movl	%r12d, %ebx
	leaq	56(%rsp), %rdi
	movq	%r9, 40(%rsp)
	leal	0(,%r8,4), %edx
	imull	%r15d, %ebx
	movslq	%edx, %rdx
	movq	%rax, 32(%rsp)
	movq	%rdx, %rsi
	movq	%rdx, 24(%rsp)
.LEHB30:
	call	_Z10cudaMallocIfE9cudaErrorPPT_m.isra.0
	movl	192(%rsp), %eax
	leaq	64(%rsp), %rdi
	imull	%r14d, %ebx
	leal	0(,%rax,4), %ebp
	movslq	%ebp, %rbp
	movslq	%ebx, %r13
	sall	$2, %ebx
	movq	%rbp, %rsi
	call	_Z10cudaMallocIfE9cudaErrorPPT_m.isra.0
	movslq	%ebx, %rbx
	leaq	72(%rsp), %rdi
	movq	%rbx, %rsi
	call	_Z10cudaMallocIfE9cudaErrorPPT_m.isra.0
	movq	32(%rsp), %rax
	movq	24(%rsp), %rdx
	movl	$1, %ecx
	movq	56(%rsp), %rdi
	movq	(%rax), %rsi
	call	cudaMemcpy@PLT
	movq	40(%rsp), %r9
	movq	64(%rsp), %rdi
	movq	%rbp, %rdx
	movl	$1, %ecx
	movq	(%r9), %rsi
	call	cudaMemcpy@PLT
	cvtsi2sdl	%r14d, %xmm0
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	addsd	.LC12(%rip), %xmm0
	subsd	.LC13(%rip), %xmm0
	movl	$64, %ecx
	mulsd	.LC14(%rip), %xmm0
	cvttsd2sil	%xmm0, %eax
	movl	%eax, 80(%rsp)
	movl	%eax, 84(%rsp)
	leal	63(%r12), %eax
	movq	80(%rsp), %rdi
	cltd
	idivl	%ecx
	movabsq	$17179869188, %rdx
	movl	$64, %ecx
	movl	%eax, %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L171
	pushq	%rax
	.cfi_def_cfa_offset 184
	movl	%r12d, %ecx
	movl	%r15d, %edx
	movl	%r14d, %edi
	pushq	80(%rsp)
	.cfi_def_cfa_offset 192
	movq	80(%rsp), %r9
	movq	72(%rsp), %r8
	movl	36(%rsp), %esi
	call	_Z36__device_stub__Z9multiplyDiiiiPfS_S_iiiiPfS_S_
	popq	%rdx
	.cfi_def_cfa_offset 184
	popq	%rcx
	.cfi_def_cfa_offset 176
.L171:
	movabsq	$2305843009213693950, %rax
	cmpq	%r13, %rax
	jb	.L172
	leaq	0(,%r13,4), %rdi
	call	_Znam@PLT
.LEHE30:
	movq	72(%rsp), %rsi
	movl	$2, %ecx
	movq	%rbx, %rdx
	movq	%rax, %rdi
	movq	8(%rsp), %rax
	movq	%rdi, (%rax)
.LEHB31:
	call	cudaMemcpy@PLT
.LEHE31:
	jmp	.L185
.L172:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L183
.LEHB32:
	call	__cxa_throw_bad_array_new_length@PLT
.LEHE32:
.L185:
	movq	56(%rsp), %rdi
.LEHB33:
	call	cudaFree@PLT
	movq	64(%rsp), %rdi
	call	cudaFree@PLT
	movq	72(%rsp), %rdi
	call	cudaFree@PLT
.LEHE33:
	jmp	.L186
.L178:
	endbr64
	movq	%rax, %rbx
.L175:
	movq	8(%rsp), %rdi
	call	_ZNSt10unique_ptrIA_fSt14default_deleteIS0_EED1Ev
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L183
	movq	%rbx, %rdi
.LEHB34:
	call	_Unwind_Resume@PLT
.LEHE34:
.L186:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L177
.L183:
	call	__stack_chk_fail@PLT
.L177:
	movq	8(%rsp), %rax
	addq	$120, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4309:
	.section	.gcc_except_table
.LLSDA4309:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE4309-.LLSDACSB4309
.LLSDACSB4309:
	.uleb128 .LEHB30-.LFB4309
	.uleb128 .LEHE30-.LEHB30
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB31-.LFB4309
	.uleb128 .LEHE31-.LEHB31
	.uleb128 .L178-.LFB4309
	.uleb128 0
	.uleb128 .LEHB32-.LFB4309
	.uleb128 .LEHE32-.LEHB32
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB33-.LFB4309
	.uleb128 .LEHE33-.LEHB33
	.uleb128 .L178-.LFB4309
	.uleb128 0
	.uleb128 .LEHB34-.LFB4309
	.uleb128 .LEHE34-.LEHB34
	.uleb128 0
	.uleb128 0
.LLSDACSE4309:
	.text
	.size	_Z12CUDAmultiplyRSt10unique_ptrIA_fSt14default_deleteIS0_EERS_IA_iS1_IS5_EEiiS4_S8_ii, .-_Z12CUDAmultiplyRSt10unique_ptrIA_fSt14default_deleteIS0_EERS_IA_iS1_IS5_EEiiS4_S8_ii
	.globl	_Z10CUDArotateRSt10unique_ptrIA_fSt14default_deleteIS0_EERS_IA_iS1_IS5_EEii
	.type	_Z10CUDArotateRSt10unique_ptrIA_fSt14default_deleteIS0_EERS_IA_iS1_IS5_EEii, @function
_Z10CUDArotateRSt10unique_ptrIA_fSt14default_deleteIS0_EERS_IA_iS1_IS5_EEii:
.LFB4310:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA4310
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movl	$1, %r15d
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movq	%rdi, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movq	%rsi, %rbx
	subq	$72, %rsp
	.cfi_def_cfa_offset 128
	movq	%fs:40, %rsi
	movq	%rsi, 56(%rsp)
	movq	(%rdx), %rsi
	movl	$2, %edx
	movl	(%rsi), %r13d
	movl	4(%rsi), %r14d
.L188:
	cmpl	%edx, %ecx
	jle	.L203
	imull	(%rsi,%rdx,4), %r15d
	incq	%rdx
	jmp	.L188
.L203:
	movslq	%r8d, %r12
	sall	$2, %r8d
	leaq	16(%rsp), %rdi
	movslq	%r8d, %rax
	movq	%rax, %rsi
	movq	%rax, 8(%rsp)
.LEHB35:
	call	_Z10cudaMallocIfE9cudaErrorPPT_m.isra.0
	movq	8(%rsp), %rsi
	leaq	24(%rsp), %rdi
	call	_Z10cudaMallocIfE9cudaErrorPPT_m.isra.0
	movq	8(%rsp), %rdx
	movq	(%rbx), %rsi
	movl	$1, %ecx
	movq	16(%rsp), %rdi
	call	cudaMemcpy@PLT
	cvtsi2sdl	%r13d, %xmm0
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	addsd	.LC12(%rip), %xmm0
	subsd	.LC13(%rip), %xmm0
	movl	$64, %ecx
	mulsd	.LC14(%rip), %xmm0
	cvttsd2sil	%xmm0, %eax
	movl	%eax, 32(%rsp)
	movl	%eax, 36(%rsp)
	leal	63(%r15), %eax
	movq	32(%rsp), %rdi
	cltd
	idivl	%ecx
	movabsq	$17179869188, %rdx
	movl	$64, %ecx
	movl	%eax, %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L190
	movq	24(%rsp), %r8
	movq	16(%rsp), %rcx
	movl	%r15d, %edx
	movl	%r14d, %esi
	movl	%r13d, %edi
	call	_Z31__device_stub__Z7rotateDiiiPfS_iiiPfS_
.L190:
	movabsq	$2305843009213693950, %rax
	cmpq	%r12, %rax
	jb	.L191
	leaq	0(,%r12,4), %rdi
	call	_Znam@PLT
.LEHE35:
	movq	8(%rsp), %rdx
	movq	24(%rsp), %rsi
	movl	$2, %ecx
	movq	%rax, 0(%rbp)
	movq	%rax, %rdi
.LEHB36:
	call	cudaMemcpy@PLT
.LEHE36:
	jmp	.L204
.L191:
	movq	56(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L202
.LEHB37:
	call	__cxa_throw_bad_array_new_length@PLT
.LEHE37:
.L204:
	movq	16(%rsp), %rdi
.LEHB38:
	call	cudaFree@PLT
	movq	24(%rsp), %rdi
	call	cudaFree@PLT
.LEHE38:
	jmp	.L205
.L197:
	endbr64
	movq	%rax, %rbx
.L194:
	movq	%rbp, %rdi
	call	_ZNSt10unique_ptrIA_fSt14default_deleteIS0_EED1Ev
	movq	56(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L202
	movq	%rbx, %rdi
.LEHB39:
	call	_Unwind_Resume@PLT
.LEHE39:
.L205:
	movq	56(%rsp), %rax
	subq	%fs:40, %rax
	je	.L196
.L202:
	call	__stack_chk_fail@PLT
.L196:
	addq	$72, %rsp
	.cfi_def_cfa_offset 56
	movq	%rbp, %rax
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4310:
	.section	.gcc_except_table
.LLSDA4310:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE4310-.LLSDACSB4310
.LLSDACSB4310:
	.uleb128 .LEHB35-.LFB4310
	.uleb128 .LEHE35-.LEHB35
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB36-.LFB4310
	.uleb128 .LEHE36-.LEHB36
	.uleb128 .L197-.LFB4310
	.uleb128 0
	.uleb128 .LEHB37-.LFB4310
	.uleb128 .LEHE37-.LEHB37
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB38-.LFB4310
	.uleb128 .LEHE38-.LEHB38
	.uleb128 .L197-.LFB4310
	.uleb128 0
	.uleb128 .LEHB39-.LFB4310
	.uleb128 .LEHE39-.LEHB39
	.uleb128 0
	.uleb128 0
.LLSDACSE4310:
	.text
	.size	_Z10CUDArotateRSt10unique_ptrIA_fSt14default_deleteIS0_EERS_IA_iS1_IS5_EEii, .-_Z10CUDArotateRSt10unique_ptrIA_fSt14default_deleteIS0_EERS_IA_iS1_IS5_EEii
	.globl	_Z14CUDAmaxPoolingRSt10unique_ptrIA_fSt14default_deleteIS0_EERS_IA_iS1_IS5_EEiiiiii
	.type	_Z14CUDAmaxPoolingRSt10unique_ptrIA_fSt14default_deleteIS0_EERS_IA_iS1_IS5_EEiiiiii, @function
_Z14CUDAmaxPoolingRSt10unique_ptrIA_fSt14default_deleteIS0_EERS_IA_iS1_IS5_EEiiiiii:
.LFB4311:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA4311
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movq	%rsi, %r15
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movl	%r9d, %r12d
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movl	$1, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movq	%rdi, %rbx
	subq	$88, %rsp
	.cfi_def_cfa_offset 144
	movq	(%rdx), %rdx
	movq	%fs:40, %rax
	movq	%rax, 72(%rsp)
	xorl	%eax, %eax
	movl	160(%rsp), %esi
	movl	4(%rdx), %eax
	movl	(%rdx), %r14d
	movl	%eax, 20(%rsp)
	movl	$2, %eax
.L207:
	cmpl	%eax, %ecx
	jle	.L222
	imull	(%rdx,%rax,4), %ebp
	incq	%rax
	jmp	.L207
.L222:
	movl	%r14d, %eax
	leaq	32(%rsp), %rdi
	subl	%r12d, %eax
	addl	152(%rsp), %eax
	cltd
	idivl	152(%rsp)
	movl	%eax, %ecx
	movl	20(%rsp), %eax
	subl	144(%rsp), %eax
	addl	%esi, %eax
	cltd
	idivl	%esi
	leal	0(,%r8,4), %edx
	movslq	%edx, %rdx
	movq	%rdx, %rsi
	movq	%rdx, 24(%rsp)
	imull	%ecx, %eax
	imull	%ebp, %eax
	movl	%eax, 8(%rsp)
	movslq	%eax, %r13
.LEHB40:
	call	_Z10cudaMallocIfE9cudaErrorPPT_m.isra.0
	movl	8(%rsp), %eax
	leaq	40(%rsp), %rdi
	sall	$2, %eax
	cltq
	movq	%rax, %rsi
	movq	%rax, 8(%rsp)
	call	_Z10cudaMallocIfE9cudaErrorPPT_m.isra.0
	movq	24(%rsp), %rdx
	movq	(%r15), %rsi
	movl	$1, %ecx
	movq	32(%rsp), %rdi
	call	cudaMemcpy@PLT
	cvtsi2sdl	%r14d, %xmm0
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	addsd	.LC12(%rip), %xmm0
	subsd	.LC13(%rip), %xmm0
	movl	$64, %ecx
	mulsd	.LC14(%rip), %xmm0
	cvttsd2sil	%xmm0, %eax
	movl	%eax, 48(%rsp)
	movl	%eax, 52(%rsp)
	leal	63(%rbp), %eax
	movq	48(%rsp), %rdi
	cltd
	idivl	%ecx
	movabsq	$17179869188, %rdx
	movl	$64, %ecx
	movl	%eax, %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L209
	pushq	%rax
	.cfi_def_cfa_offset 152
	movl	%r12d, %r8d
	movl	%ebp, %edx
	movl	%r14d, %edi
	pushq	48(%rsp)
	.cfi_def_cfa_offset 160
	pushq	48(%rsp)
	.cfi_def_cfa_offset 168
	movl	176(%rsp), %eax
	pushq	%rax
	.cfi_def_cfa_offset 176
	movl	184(%rsp), %r9d
	movl	176(%rsp), %ecx
	movl	52(%rsp), %esi
	call	_Z40__device_stub__Z11maxPoolingDiiiiiiiPfS_iiiiiiiPfS_
	addq	$32, %rsp
	.cfi_def_cfa_offset 144
.L209:
	movabsq	$2305843009213693950, %rax
	cmpq	%r13, %rax
	jb	.L210
	leaq	0(,%r13,4), %rdi
	call	_Znam@PLT
.LEHE40:
	movq	8(%rsp), %rdx
	movq	40(%rsp), %rsi
	movl	$2, %ecx
	movq	%rax, (%rbx)
	movq	%rax, %rdi
.LEHB41:
	call	cudaMemcpy@PLT
.LEHE41:
	jmp	.L223
.L210:
	movq	72(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L221
.LEHB42:
	call	__cxa_throw_bad_array_new_length@PLT
.LEHE42:
.L223:
	movq	32(%rsp), %rdi
.LEHB43:
	call	cudaFree@PLT
	movq	40(%rsp), %rdi
	call	cudaFree@PLT
.LEHE43:
	jmp	.L224
.L216:
	endbr64
	movq	%rax, %rbp
.L213:
	movq	%rbx, %rdi
	call	_ZNSt10unique_ptrIA_fSt14default_deleteIS0_EED1Ev
	movq	72(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L221
	movq	%rbp, %rdi
.LEHB44:
	call	_Unwind_Resume@PLT
.LEHE44:
.L224:
	movq	72(%rsp), %rax
	subq	%fs:40, %rax
	je	.L215
.L221:
	call	__stack_chk_fail@PLT
.L215:
	addq	$88, %rsp
	.cfi_def_cfa_offset 56
	movq	%rbx, %rax
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4311:
	.section	.gcc_except_table
.LLSDA4311:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE4311-.LLSDACSB4311
.LLSDACSB4311:
	.uleb128 .LEHB40-.LFB4311
	.uleb128 .LEHE40-.LEHB40
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB41-.LFB4311
	.uleb128 .LEHE41-.LEHB41
	.uleb128 .L216-.LFB4311
	.uleb128 0
	.uleb128 .LEHB42-.LFB4311
	.uleb128 .LEHE42-.LEHB42
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB43-.LFB4311
	.uleb128 .LEHE43-.LEHB43
	.uleb128 .L216-.LFB4311
	.uleb128 0
	.uleb128 .LEHB44-.LFB4311
	.uleb128 .LEHE44-.LEHB44
	.uleb128 0
	.uleb128 0
.LLSDACSE4311:
	.text
	.size	_Z14CUDAmaxPoolingRSt10unique_ptrIA_fSt14default_deleteIS0_EERS_IA_iS1_IS5_EEiiiiii, .-_Z14CUDAmaxPoolingRSt10unique_ptrIA_fSt14default_deleteIS0_EERS_IA_iS1_IS5_EEiiiiii
	.globl	_Z16CUDApoolingDerivRSt10unique_ptrIA_fSt14default_deleteIS0_EERS_IA_iS1_IS5_EEiiS4_S8_iiiiii
	.type	_Z16CUDApoolingDerivRSt10unique_ptrIA_fSt14default_deleteIS0_EERS_IA_iS1_IS5_EEiiS4_S8_iiiiii, @function
_Z16CUDApoolingDerivRSt10unique_ptrIA_fSt14default_deleteIS0_EERS_IA_iS1_IS5_EEiiS4_S8_iiiiii:
.LFB4312:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA4312
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	movq	%rsi, %r14
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movq	%rdi, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movl	$1, %ebx
	subq	$104, %rsp
	.cfi_def_cfa_offset 160
	movq	%fs:40, %rsi
	movq	%rsi, 88(%rsp)
	movq	(%rdx), %rsi
	movl	$2, %edx
	movl	4(%rsi), %eax
	movl	(%rsi), %r13d
	movl	%eax, 20(%rsp)
.L226:
	cmpl	%edx, %ecx
	jle	.L241
	imull	(%rsi,%rdx,4), %ebx
	incq	%rdx
	jmp	.L226
.L241:
	movslq	%r8d, %r12
	sall	$2, %r8d
	leaq	40(%rsp), %rdi
	movq	%r9, 24(%rsp)
	movslq	%r8d, %rax
	movq	%rax, %rsi
	movq	%rax, 8(%rsp)
.LEHB45:
	call	_Z10cudaMallocIfE9cudaErrorPPT_m.isra.0
	movl	176(%rsp), %eax
	leaq	48(%rsp), %rdi
	leal	0(,%rax,4), %r15d
	movslq	%r15d, %r15
	movq	%r15, %rsi
	call	_Z10cudaMallocIfE9cudaErrorPPT_m.isra.0
	movq	8(%rsp), %rsi
	leaq	56(%rsp), %rdi
	call	_Z10cudaMallocIfE9cudaErrorPPT_m.isra.0
	movq	8(%rsp), %rdx
	movq	(%r14), %rsi
	movl	$1, %ecx
	movq	40(%rsp), %rdi
	call	cudaMemcpy@PLT
	movq	24(%rsp), %r9
	movq	48(%rsp), %rdi
	movq	%r15, %rdx
	movl	$1, %ecx
	movq	(%r9), %rsi
	call	cudaMemcpy@PLT
	cvtsi2sdl	%r13d, %xmm0
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	addsd	.LC12(%rip), %xmm0
	subsd	.LC13(%rip), %xmm0
	movl	$64, %ecx
	mulsd	.LC14(%rip), %xmm0
	cvttsd2sil	%xmm0, %eax
	movl	%eax, 64(%rsp)
	movl	%eax, 68(%rsp)
	leal	63(%rbx), %eax
	movq	64(%rsp), %rdi
	cltd
	idivl	%ecx
	movabsq	$17179869188, %rdx
	movl	$64, %ecx
	movl	%eax, %esi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L228
	pushq	56(%rsp)
	.cfi_def_cfa_offset 168
	movl	%ebx, %edx
	movl	%r13d, %edi
	pushq	56(%rsp)
	.cfi_def_cfa_offset 176
	pushq	56(%rsp)
	.cfi_def_cfa_offset 184
	movl	232(%rsp), %eax
	pushq	%rax
	.cfi_def_cfa_offset 192
	movl	232(%rsp), %r9d
	movl	224(%rsp), %r8d
	movl	216(%rsp), %ecx
	movl	52(%rsp), %esi
	call	_Z44__device_stub__Z13poolingDerivDiiiiiiiPfS_S_iiiiiiiPfS_S_
	addq	$32, %rsp
	.cfi_def_cfa_offset 160
.L228:
	movabsq	$2305843009213693950, %rax
	cmpq	%r12, %rax
	jb	.L229
	leaq	0(,%r12,4), %rdi
	call	_Znam@PLT
.LEHE45:
	movq	8(%rsp), %rdx
	movq	56(%rsp), %rsi
	movl	$2, %ecx
	movq	%rax, 0(%rbp)
	movq	%rax, %rdi
.LEHB46:
	call	cudaMemcpy@PLT
.LEHE46:
	jmp	.L242
.L229:
	movq	88(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L240
.LEHB47:
	call	__cxa_throw_bad_array_new_length@PLT
.LEHE47:
.L242:
	movq	40(%rsp), %rdi
.LEHB48:
	call	cudaFree@PLT
	movq	48(%rsp), %rdi
	call	cudaFree@PLT
	movq	56(%rsp), %rdi
	call	cudaFree@PLT
.LEHE48:
	jmp	.L243
.L235:
	endbr64
	movq	%rax, %rbx
.L232:
	movq	%rbp, %rdi
	call	_ZNSt10unique_ptrIA_fSt14default_deleteIS0_EED1Ev
	movq	88(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L240
	movq	%rbx, %rdi
.LEHB49:
	call	_Unwind_Resume@PLT
.LEHE49:
.L243:
	movq	88(%rsp), %rax
	subq	%fs:40, %rax
	je	.L234
.L240:
	call	__stack_chk_fail@PLT
.L234:
	addq	$104, %rsp
	.cfi_def_cfa_offset 56
	movq	%rbp, %rax
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4312:
	.section	.gcc_except_table
.LLSDA4312:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE4312-.LLSDACSB4312
.LLSDACSB4312:
	.uleb128 .LEHB45-.LFB4312
	.uleb128 .LEHE45-.LEHB45
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB46-.LFB4312
	.uleb128 .LEHE46-.LEHB46
	.uleb128 .L235-.LFB4312
	.uleb128 0
	.uleb128 .LEHB47-.LFB4312
	.uleb128 .LEHE47-.LEHB47
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB48-.LFB4312
	.uleb128 .LEHE48-.LEHB48
	.uleb128 .L235-.LFB4312
	.uleb128 0
	.uleb128 .LEHB49-.LFB4312
	.uleb128 .LEHE49-.LEHB49
	.uleb128 0
	.uleb128 0
.LLSDACSE4312:
	.text
	.size	_Z16CUDApoolingDerivRSt10unique_ptrIA_fSt14default_deleteIS0_EERS_IA_iS1_IS5_EEiiS4_S8_iiiiii, .-_Z16CUDApoolingDerivRSt10unique_ptrIA_fSt14default_deleteIS0_EERS_IA_iS1_IS5_EEiiS4_S8_iiiiii
	.globl	_Z8CUDAdupeRSt10unique_ptrIA_fSt14default_deleteIS0_EERS_IA_iS1_IS5_EEiii
	.type	_Z8CUDAdupeRSt10unique_ptrIA_fSt14default_deleteIS0_EERS_IA_iS1_IS5_EEiii, @function
_Z8CUDAdupeRSt10unique_ptrIA_fSt14default_deleteIS0_EERS_IA_iS1_IS5_EEiii:
.LFB4313:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA4313
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movq	%rsi, %rax
	movl	$1, %r10d
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	movq	%rdi, %r13
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movl	%r9d, %r12d
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	subq	$72, %rsp
	.cfi_def_cfa_offset 128
	movq	%fs:40, %rsi
	movq	%rsi, 56(%rsp)
	movq	(%rdx), %rsi
	movl	$2, %edx
	movl	(%rsi), %r15d
	movl	4(%rsi), %r11d
.L245:
	cmpl	%edx, %ecx
	jle	.L260
	imull	(%rsi,%rdx,4), %r10d
	incq	%rdx
	jmp	.L245
.L260:
	movl	%r15d, %ebx
	imull	%r10d, %r12d
	leaq	16(%rsp), %rdi
	movl	%r11d, (%rsp)
	imull	%r11d, %ebx
	leal	0(,%r8,4), %ebp
	movl	%r10d, 4(%rsp)
	movslq	%ebp, %rbp
	movq	%rax, 8(%rsp)
	movq	%rbp, %rsi
.LEHB50:
	call	_Z10cudaMallocIfE9cudaErrorPPT_m.isra.0
	imull	%r12d, %ebx
	leaq	24(%rsp), %rdi
	movslq	%ebx, %r14
	sall	$2, %ebx
	movslq	%ebx, %rbx
	movq	%rbx, %rsi
	call	_Z10cudaMallocIfE9cudaErrorPPT_m.isra.0
	movq	8(%rsp), %rax
	movq	16(%rsp), %rdi
	movq	%rbp, %rdx
	movl	$1, %ecx
	movq	(%rax), %rsi
	call	cudaMemcpy@PLT
	cvtsi2sdl	%r15d, %xmm0
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	addsd	.LC12(%rip), %xmm0
	subsd	.LC13(%rip), %xmm0
	movl	$64, %ecx
	mulsd	.LC14(%rip), %xmm0
	cvttsd2sil	%xmm0, %eax
	movl	%eax, 32(%rsp)
	movl	%eax, 36(%rsp)
	leal	63(%r12), %eax
	movq	32(%rsp), %rdi
	cltd
	idivl	%ecx
	movabsq	$17179869188, %rdx
	movl	$64, %ecx
	movl	%eax, %esi
	call	__cudaPushCallConfiguration@PLT
	movl	(%rsp), %r11d
	movl	4(%rsp), %r10d
	testl	%eax, %eax
	jne	.L247
	movq	24(%rsp), %r9
	movq	16(%rsp), %r8
	movl	%r12d, %ecx
	movl	%r10d, %edx
	movl	%r11d, %esi
	movl	%r15d, %edi
	call	_Z30__device_stub__Z5dupeDiiiiPfS_iiiiPfS_
.L247:
	movabsq	$2305843009213693950, %rax
	cmpq	%r14, %rax
	jb	.L248
	leaq	0(,%r14,4), %rdi
	call	_Znam@PLT
.LEHE50:
	movq	24(%rsp), %rsi
	movl	$2, %ecx
	movq	%rbx, %rdx
	movq	%rax, 0(%r13)
	movq	%rax, %rdi
.LEHB51:
	call	cudaMemcpy@PLT
.LEHE51:
	jmp	.L261
.L248:
	movq	56(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L259
.LEHB52:
	call	__cxa_throw_bad_array_new_length@PLT
.LEHE52:
.L261:
	movq	16(%rsp), %rdi
.LEHB53:
	call	cudaFree@PLT
	movq	24(%rsp), %rdi
	call	cudaFree@PLT
.LEHE53:
	jmp	.L262
.L254:
	endbr64
	movq	%rax, %rbx
.L251:
	movq	%r13, %rdi
	call	_ZNSt10unique_ptrIA_fSt14default_deleteIS0_EED1Ev
	movq	56(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L259
	movq	%rbx, %rdi
.LEHB54:
	call	_Unwind_Resume@PLT
.LEHE54:
.L262:
	movq	56(%rsp), %rax
	subq	%fs:40, %rax
	je	.L253
.L259:
	call	__stack_chk_fail@PLT
.L253:
	addq	$72, %rsp
	.cfi_def_cfa_offset 56
	movq	%r13, %rax
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4313:
	.section	.gcc_except_table
.LLSDA4313:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE4313-.LLSDACSB4313
.LLSDACSB4313:
	.uleb128 .LEHB50-.LFB4313
	.uleb128 .LEHE50-.LEHB50
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB51-.LFB4313
	.uleb128 .LEHE51-.LEHB51
	.uleb128 .L254-.LFB4313
	.uleb128 0
	.uleb128 .LEHB52-.LFB4313
	.uleb128 .LEHE52-.LEHB52
	.uleb128 0
	.uleb128 0
	.uleb128 .LEHB53-.LFB4313
	.uleb128 .LEHE53-.LEHB53
	.uleb128 .L254-.LFB4313
	.uleb128 0
	.uleb128 .LEHB54-.LFB4313
	.uleb128 .LEHE54-.LEHB54
	.uleb128 0
	.uleb128 0
.LLSDACSE4313:
	.text
	.size	_Z8CUDAdupeRSt10unique_ptrIA_fSt14default_deleteIS0_EERS_IA_iS1_IS5_EEiii, .-_Z8CUDAdupeRSt10unique_ptrIA_fSt14default_deleteIS0_EERS_IA_iS1_IS5_EEiii
	.section	.nvFatBinSegment,"aw"
	.align 8
	.type	_ZL15__fatDeviceText, @object
	.size	_ZL15__fatDeviceText, 24
_ZL15__fatDeviceText:
	.long	1180844977
	.long	1
	.quad	fatbinData
	.quad	0
	.local	_ZL20__cudaFatCubinHandle
	.comm	_ZL20__cudaFatCubinHandle,8,8
	.section	.rodata.cst8,"aM",@progbits,8
	.align 8
.LC12:
	.long	0
	.long	1074790400
	.align 8
.LC13:
	.long	0
	.long	1072693248
	.align 8
.LC14:
	.long	0
	.long	1070596096
	.hidden	DW.ref.__gxx_personality_v0
	.weak	DW.ref.__gxx_personality_v0
	.section	.data.rel.local.DW.ref.__gxx_personality_v0,"awG",@progbits,DW.ref.__gxx_personality_v0,comdat
	.align 8
	.type	DW.ref.__gxx_personality_v0, @object
	.size	DW.ref.__gxx_personality_v0, 8
DW.ref.__gxx_personality_v0:
	.quad	__gxx_personality_v0
	.ident	"GCC: (Ubuntu 13.3.0-6ubuntu2~24.04) 13.3.0"
	.section	.note.GNU-stack,"",@progbits
	.section	.note.gnu.property,"a"
	.align 8
	.long	1f - 0f
	.long	4f - 1f
	.long	5
0:
	.string	"GNU"
1:
	.align 8
	.long	0xc0000002
	.long	3f - 2f
2:
	.long	0x3
3:
	.align 8
4:

#pragma once
#include <hip/hip_runtime.h>


// Perform the error checking on the python higher level part
// Maybe we want to change the block sizes so that there are more threads for different GPU's
// Why does it have "in_ptr" followed by "ptr_size" .... ? - reformat this

#include <iostream>
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(hipError_t code, const char *file, int line, bool abort=true) {
    if (code != hipSuccess) {
        std::cout << hipGetErrorString(code) << " " << file << " " << line << std::endl;
        if (abort) exit(code);
    }
}

#include <memory>
#include <cmath>

// Perform cleanup and standardization of different function
// Organize functions into classes
// Nothing has been tested in more than four dimensions...?

std::unique_ptr<float[]> CUDAaddElementwise(std::unique_ptr<float[]>& in_ptr1, std::unique_ptr<float[]>& in_ptr2, int ptr_size);
std::unique_ptr<float[]> CUDAmultiplyElementwise(std::unique_ptr<float[]>& in_ptr1, std::unique_ptr<float[]>& in_ptr2, int ptr_size);
std::unique_ptr<float[]> CUDAdivideElementwise(std::unique_ptr<float[]>& in_ptr1, std::unique_ptr<float[]>& in_ptr2, int ptr_size);
std::unique_ptr<float[]> CUDAsubtractElementwise(std::unique_ptr<float[]>& in_ptr1, std::unique_ptr<float[]>& in_ptr2, int ptr_size);
std::unique_ptr<float[]> CUDApowerElementwise(std::unique_ptr<float[]>& in_ptr1, std::unique_ptr<float[]>& in_ptr2, int ptr_size);

std::unique_ptr<float[]> CUDAtranspose(std::unique_ptr<float[]>& in_ptr1, std::unique_ptr<int[]>& in_ptr1_dims, int in_ptr1_dims_size, int ptr1_size);
std::unique_ptr<float[]> CUDAmultiply(std::unique_ptr<float[]>& in_ptr1, std::unique_ptr<int[]>& in_ptr1_dims, int in_ptr1_dims_size, int ptr1_size, std::unique_ptr<float[]>& in_ptr2, std::unique_ptr<int[]>& in_ptr2_dims, int in_ptr2_dims_size, int ptr2_size);

// How am I going to reverse this pooling layer
// Pooling layer for back propogation
std::unique_ptr<float[]> CUDApoolingDeriv(std::unique_ptr<float[]>& in_ptr1, std::unique_ptr<int[]>& in_ptr1_dims, int in_ptr1_dims_size, int ptr1_size, std::unique_ptr<float[]>& in_ptr2, std::unique_ptr<int[]>& in_ptr2_dims, int in_ptr2_dims_size, int ptr2_size,  int kernel_cols, int kernel_rows, int stride_cols, int stride_rows);
std::unique_ptr<float[]> CUDAmaxPooling(std::unique_ptr<float[]>& in_ptr1, std::unique_ptr<int[]>& in_ptr1_dims, int in_ptr1_dims_size, int ptr1_size, int kernel_cols, int kernel_rows, int stride_cols, int stride_rows);

std::unique_ptr<float[]> CUDApad(std::unique_ptr<float[]>& in_ptr1, std::unique_ptr<int[]>& in_ptr1_dims, int in_ptr1_dims_size, int ptr1_size, int pad_left, int pad_right, int pad_up, int pad_down, int pad_between_cols, int pad_between_rows);

std::unique_ptr<float[]> CUDArotate(std::unique_ptr<float[]>& in_ptr1, std::unique_ptr<int[]>& in_ptr1_dims, int in_ptr1_dims_size, int ptr1_size);

std::unique_ptr<float[]> CUDAdupe(std::unique_ptr<float[]>& in_ptr1, std::unique_ptr<int[]>& in_ptr1_dims, int in_ptr1_dims_size, int ptr1_size, int dupe_size);

// What if for the convolutional tensor I simply dont do the compilation sum which would make it tricky to backprop through
// Might have to expand this out to the fourth dimension for true parallel processing
// Dont forget the bias term! --- This can be done with a simple addition afterwards --- Training for this might be a bit painful
std::unique_ptr<float[]> CUDAconvolution(std::unique_ptr<float[]>& in_ptr1, std::unique_ptr<int[]>& in_ptr1_dims, int in_ptr1_dims_size, int ptr1_size, std::unique_ptr<float[]>& in_ptr2, std::unique_ptr<int[]>& in_ptr2_dims, int in_ptr2_dims_size, int ptr2_size, int stride_cols, int stride_rows);
	.text
	.file	"tensoroperations.hip"
                                        # Start of file scope inline assembly
	.globl	_ZSt21ios_base_library_initv

                                        # End of file scope inline assembly
	.type	__hip_cuid_,@object             # @__hip_cuid_
	.bss
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               # 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_

	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.amdgpu_metadata
---
amdhsa.kernels:  []
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
