<commit_before>// The GEMM/im2col implementation used here is:
// - adapted from clnn: https://github.com/hughperkins/clnn/blob/master/lib/THCLNN/im2col.cpp), which was
// - ported from cunn: https://github.com/torch/cunn/blob/master/lib/THCUNN/im2col.h), which was
// - adapted from Caffe: https://github.com/BVLC/caffe/blob/master/src/caffe/layers/conv_layer.cu)

// other convolutional implementatinos can be added in the future, but GEMM/im2col gives a solid works-everywhere base
// to start from

#include "cocl_dnn_gemm.h"

#include "cocl/cocl.h"
#include "cocl/cocl_dnn.h"
#include "cocl/cocl_memory.h"
// #include "cocl/cocl_blas.h"
#include <clblast_c.h>

#include <iostream>
using namespace std;

namespace cocl {
namespace dnn {
namespace gemm_im2col {

static string get_im2col_sourcecode();
static string get_col2im_sourcecode();

CoclDnnGeometryType getColumnsNumElements(
        cudnnHandle_t handle,
        cudnnTensorDescriptor_t srcTensor,
        cudnnFilterDescriptor_t filter,
        cudnnConvolutionDescriptor_t conv,
        cudnnTensorDescriptor_t dstTensor) {

    CoclDnnGeometryType outN, outC, outH, outW;
    cudnnGetConvolution2dForwardOutputDim(
        conv,
        srcTensor,
        filter,
        &outN, &outC, &outH, &outW);

    CoclDnnGeometryType inC = srcTensor->C;
    CoclDnnGeometryType rows = inC * filter->kW * filter->kH;
    CoclDnnGeometryType cols = outH * outW;
    return rows * cols;
}

// void im2col(THClTensor* im, const int channels,
//         const int height, const int width, const int ksize_h, const int ksize_w, const int pad_h,
//         const int pad_w, const int stride_h, const int stride_w, THClTensor* col) {
//     // We are going to launch channels * height_col * width_col kernels, each
//     // kernel responsible for copying a single-channel grid.
//     int height_col = (height + 2 * pad_h - ksize_h) / stride_h + 1;
//     int width_col = (width + 2 * pad_w - ksize_w) / stride_w + 1;
//     int num_kernels = channels * height_col * width_col;

    // CLKernel *getKernelForNameCl("im2col_kernel", im2col_sourcecode);

    // THClKernels k(state, kernel);
    // k.in(num_kernels);
    // k.in(im);
    // k.in(height);
    // k.in(width);
    // k.in(ksize_h);
    // k.in(ksize_w);
    // k.in(pad_h);
    // k.in(pad_w);
    // k.in(stride_h);
    // k.in(stride_w);
    // k.in(height_col);
    // k.in(width_col);
    // k.out(col);

    // k.run(GET_BLOCKS(state, num_kernels), getNumThreads(state));
// }

// void col2im(THClTensor* col, const int channels,
//         const int height, const int width, const int patch_h, const int patch_w, const int pad_h,
//         const int pad_w, const int stride_h, const int stride_w, THClTensor* im) {
//     int height_col = (height + 2 * pad_h - patch_h) / stride_h + 1;
//     int width_col = (width + 2 * pad_w - patch_w) / stride_w + 1;
//     int num_kernels = channels * height * width;
//     // To avoid involving atomic operations, we will launch one kernel per
//     // bottom dimension, and then in the kernel add up the top dimensions.

//     EasyCL *cl = im->storage->cl;
//     std::string uniqueName = "SpatialConvolutionMM::col2im";
//     CLKernel *kernel = 0;
//     if(cl->kernelExists(uniqueName)) {
//         kernel = cl->getKernel(uniqueName);
//     } else {
//         TemplatedKernel kernelBuilder(cl);
//         kernel = kernelBuilder.buildKernel(uniqueName, "SpatialConvolutionMM.cl",
//             SpatialConvolutionMM_getKernelTemplate(), "col2im_kernel");
//     }

//     THClKernels k(state, kernel);
//     k.in(num_kernels);
//     k.in(col);
//     k.in(height);
//     k.in(width);
//     k.in(channels);

//     k.in(patch_h);
//     k.in(patch_w);
//     k.in(pad_h);<commit_msg>start to add im2col<commit_after>// The GEMM/im2col implementation used here is:
// - adapted from clnn: https://github.com/hughperkins/clnn/blob/master/lib/THCLNN/im2col.cpp), which was
// - ported from cunn: https://github.com/torch/cunn/blob/master/lib/THCUNN/im2col.h), which was
// - adapted from Caffe: https://github.com/BVLC/caffe/blob/master/src/caffe/layers/conv_layer.cu)

// other convolutional implementatinos can be added in the future, but GEMM/im2col gives a solid works-everywhere base
// to start from

#include "cocl_dnn_gemm.h"

#include "cocl/cocl.h"
#include "cocl/cocl_dnn.h"
#include "cocl/cocl_memory.h"
#include "cocl/hostside_opencl_funcs.h"
// #include "cocl/cocl_blas.h"
#include <clblast_c.h>

#include <iostream>
using namespace std;

namespace cocl {
namespace dnn {
namespace gemm_im2col {

static string get_im2col_sourcecode();
static string get_col2im_sourcecode();

CoclDnnGeometryType getColumnsNumElements(
        cudnnHandle_t handle,
        cudnnTensorDescriptor_t srcTensor,
        cudnnFilterDescriptor_t filter,
        cudnnConvolutionDescriptor_t conv,
        cudnnTensorDescriptor_t dstTensor) {

    CoclDnnGeometryType outN, outC, outH, outW;
    cudnnGetConvolution2dForwardOutputDim(
        conv,
        srcTensor,
        filter,
        &outN, &outC, &outH, &outW);

    CoclDnnGeometryType inC = srcTensor->C;
    CoclDnnGeometryType rows = inC * filter->kW * filter->kH;
    CoclDnnGeometryType cols = outH * outW;
    return rows * cols;
}

inline int getNumThreads() {
  // int blockSize = 1024;
  // int maxWorkgroupSize = ((easycl::DeviceInfo *)state->deviceInfoByDevice[state->currentDevice])->maxWorkGroupSize;
  // if( blockSize > maxWorkgroupSize ) {
  //   blockSize = maxWorkgroupSize;
  // }
  // return blockSize;
    return 256; // just hardcode to 256 for now, which covers amd, intel, nvidia, just not always most efficiently, but
                // kind of ok
}

// CL: number of blocks for threads.
inline int GET_BLOCKS(const int N) {
  return (N + getNumThreads() - 1) / getNumThreads();
}

void im2col(cl_mem im_buf, size_t im_offset, const CoclDnnGeometryType channels,
        const CoclDnnGeometryType height,
        const CoclDnnGeometryType width,
        const CoclDnnGeometryType ksize_h,
        const CoclDnnGeometryType ksize_w,
        const CoclDnnGeometryType pad_h,
        const CoclDnnGeometryType pad_w,
        const CoclDnnGeometryType stride_h,
        const CoclDnnGeometryType stride_w,
        cl_mem col_buf, size_t col_offset
        ) {
    // We are going to launch channels * height_col * width_col kernels, each
    // kernel responsible for copying a single-channel grid.
    int height_col = (height + 2 * pad_h - ksize_h) / stride_h + 1;
    int width_col = (width + 2 * pad_w - ksize_w) / stride_w + 1;
    int num_kernels = channels * height_col * width_col;

    easycl::CLKernel *kernel = getKernelForNameCl("im2col_kernel", get_im2col_sourcecode());

    kernel->in(num_kernels);
    kernel->inout(&im_buf);
    kernel->in(im_offset);
    kernel->in(height);
    kernel->in(width);
    kernel->in(ksize_h);
    kernel->in(ksize_w);
    kernel->in(pad_h);
    kernel->in(pad_w);
    kernel->in(stride_h);
    kernel->in(stride_w);
    kernel->in(height_col);
    kernel->in(width_col);
    kernel->inout(&col_buf);
    kernel->in(col_offset);
    kernel->run_1d(GET_BLOCKS(num_kernels), getNumThreads());
}

// void col2im(THClTensor* col, const int channels,
//         const int height, const int width, const int patch_h, const int patch_w, const int pad_h,
//         const int pad_w, const int stride_h, const int stride_w, THClTensor* im) {
//     int height_col = (height + 2 * pad_h - patch_h) / stride_h + 1;
//     int width_col = (width + 2 * pad_w - patch_w) / stride_w + 1;
//     int num_kernels = channels * height * width;
//     // To avoid involving atomic operations, we will launch one kernel per
//     // bottom dimension, and then in the kernel add up the top dimensions.

//     EasyCL *cl = im->storage->cl;
//     std::string uniqueName = "SpatialConvolutionMM::col2im";
//     CLKernel *kernel = 0;
//     if(cl->kernelExists(uniqueName)) {
//         kernel = cl->getKernel(uniqueName);
//     } else {
//         TemplatedKernel kernelBuilder(cl);
//         kernel = kernelBuilder.buildKernel(uniqueName, "SpatialConvolutionMM.cl",
//             SpatialConvolutionMM_getKernelTemplate(), "col2im_kernel");
//     }

//     THClKernels k(state, kernel);
//     k.in(num_kernels);
//     k.in(col);
//     k.in(height);
//     k.in(width);
//     k.in(channels);

//     k.in(patch_h);
//     k.in(patch_w);
//     k.in(pad_h);
{'fraction_non_alphanumeric': 0.12689753320683111, 'fraction_numerical': 0.006997153700189753, 'mean_word_length': 2.9535864978902953, 'pattern_counts': {'":': 0, '<': 7, '<?xml version=': 0, '>': 41, 'https://': 6, 'lorem ipsum': 0, 'www.': 0, 'xml': 0}, 'pii_count': 0, 'substrings_counts': 0, 'word_list_counts': {'cursed_substrings.json': 0, 'profanity_word_list.json': 0, 'sexual_word_list.json': 0, 'zh_pornsignals.json': 0}}
{'dir': 'git-commits-cleaned', 'id': '1740066', 'n_tokens_mistral': 3025, 'n_tokens_neox': 2934, 'n_words': 682}