diff --git a/ggml/src/ggml-cpu/ggml-cpu.c b/ggml/src/ggml-cpu/ggml-cpu.c index 0d5d3a3440..16d9f0204a 100644 --- a/ggml/src/ggml-cpu/ggml-cpu.c +++ b/ggml/src/ggml-cpu/ggml-cpu.c @@ -1880,6 +1880,10 @@ static void ggml_compute_forward(struct ggml_compute_params * params, struct ggm { ggml_compute_forward_conv_2d(params, tensor); } break; + case GGML_OP_CONV_2D_IMPLICIT: + { + ggml_compute_forward_conv_2d(params, tensor); + } break; case GGML_OP_CONV_3D: { ggml_compute_forward_conv_3d(params, tensor); @@ -2256,6 +2260,7 @@ static int ggml_get_n_tasks(struct ggml_tensor * node, int n_threads) { case GGML_OP_IM2COL: case GGML_OP_IM2COL_BACK: case GGML_OP_CONV_2D: + case GGML_OP_CONV_2D_IMPLICIT: case GGML_OP_CONV_3D: case GGML_OP_CONV_2D_DW: case GGML_OP_CONV_TRANSPOSE_1D: @@ -2778,6 +2783,7 @@ struct ggml_cplan ggml_graph_plan( } } break; case GGML_OP_CONV_2D: + case GGML_OP_CONV_2D_IMPLICIT: case GGML_OP_CONV_3D: { cur = GGML_IM2COL_WORK_SIZE; diff --git a/ggml/src/ggml-cuda/conv2d-implicit.cu b/ggml/src/ggml-cuda/conv2d-implicit.cu index d1b1dc7d3c..72f8d30baf 100644 --- a/ggml/src/ggml-cuda/conv2d-implicit.cu +++ b/ggml/src/ggml-cuda/conv2d-implicit.cu @@ -1,81 +1,33 @@ #include "conv2d-implicit.cuh" #include "convert.cuh" -struct conv_params { - const int64_t IW, IH; - const int64_t OW, OH; - const int64_t KW, KH; - const int64_t ST_X, ST_Y; - const int64_t PD_X, PD_Y; - const int64_t DL_X, DL_Y; - const int64_t IC, OC; - const int64_t B; - const int64_t TOTAL; -}; +typedef struct{ + unsigned int n; //batch szie + unsigned int c; //channel number + unsigned int h; //height + unsigned int w; //width + unsigned int k; //number of filters + unsigned int r; //filter height + unsigned int s; //filter width + unsigned int u; //stride height + unsigned int v; //stride width + unsigned int p; //padding height + unsigned int q; //padding width + unsigned int d_h; //dilation height + unsigned int d_w; //dilation width + unsigned int Oh; //output height + unsigned int Ow; //output width +} param_t; -struct kernel_bounds { - int64_t y_min, y_max; - int64_t x_min, x_max; -}; -__device__ __forceinline__ int64_t max64(int64_t a, int64_t b) { - return (a > b) ? a : b; -} -__device__ __forceinline__ int64_t min64(int64_t a, int64_t b) { - return (a < b) ? a : b; -} - -__device__ __forceinline__ kernel_bounds calculate_kernel_bounds(int64_t out_x, int64_t out_y, const conv_params & P) { - kernel_bounds bounds; - bounds.y_min = max64(0, (P.PD_Y - out_y * P.ST_Y + P.DL_Y - 1) / P.DL_Y); - bounds.y_max = min64(P.KH, (P.IH + P.PD_Y - out_y * P.ST_Y + P.DL_Y - 1) / P.DL_Y); - bounds.x_min = max64(0, (P.PD_X - out_x * P.ST_X + P.DL_X - 1) / P.DL_X); - bounds.x_max = min64(P.KW, (P.IW + P.PD_X - out_x * P.ST_X + P.DL_X - 1) / P.DL_X); - return bounds; -} - -__device__ __forceinline__ int calculate_input_coord(int64_t out_coord, - int64_t kern_coord, - int64_t stride, - int64_t dilation, - int64_t padding) { - return out_coord * stride + kern_coord * dilation - padding; -} - -struct whcn_layout { - __device__ static int64_t input_index(int64_t n, int64_t c, int64_t y, int64_t x, const conv_params & P) { - return n * (P.IC * P.IW * P.IH) + c * P.IW * P.IH + y * P.IW + x; - } - - __device__ static int64_t kernel_index(int64_t c_out, int64_t c_in, int64_t ky, int64_t kx, const conv_params & P) { - return c_out * (P.IC * P.KH * P.KW) + c_in * (P.KH * P.KW) + ky * P.KW + kx; - } - - __device__ static int64_t output_index(int64_t n, int64_t c, int64_t y, int64_t x, const conv_params & P) { - return n * (P.OC * P.OW * P.OH) + c * P.OW * P.OH + y * P.OW + x; - } - - __device__ static void unpack_indices(int64_t global_idx, - const conv_params & P, - int64_t & n, - int64_t & c, - int64_t & out_y, - int64_t & out_x) { - out_x = global_idx % P.OW; - out_y = (global_idx / P.OW) % P.OH; - c = (global_idx / (P.OW * P.OH)) % P.OC; - n = global_idx / (P.OW * P.OH * P.OC); - } -}; - -template +template static __global__ void conv2d_implicit_kernel(const float * __restrict__ input, const T * __restrict__ kernel, float * __restrict__ output, - const conv_params P) { + const param_t ¶m) { - __shared__ __align__(16 * 1024) char smem[24 * 1024]; + extern __shared__ __align__(16 * 1024) char smem[]; T *smemweight = reinterpret_cast(smem); float *smeminput = reinterpret_cast(smem + 16 * 1024); @@ -151,8 +103,8 @@ static __global__ void conv2d_implicit_kernel(const float * __restrict__ input, #pragma unroll for (int i = 0; i < 4; ++i) { - int curH = posh_ori[i] + curR; // input h - int curW = posw_ori[i] + curS; // input w + int curH = posh_ori[i] + curR * param.d_h; // input h + int curW = posw_ori[i] + curS * param.d_w; // input w int inOffsetTmp = curC * inChannelOffset + curH * param.w + curW; if (curH >= 0 && curW >= 0 && curW < param.w && curH < param.h) { @@ -210,8 +162,8 @@ static __global__ void conv2d_implicit_kernel(const float * __restrict__ input, #pragma unroll for (int i = 0; i < 4; ++i) { - int curH = posh_ori[i] + curR; // input h - int curW = posw_ori[i] + curS; // input w + int curH = posh_ori[i] + curR * param.d_h; // input h + int curW = posw_ori[i] + curS * param.d_w; // input w int inOffsetTmp = curC * inChannelOffset + curH * param.w + curW; if (curH >= 0 && curW >= 0 && curW < param.w && curH < param.h) { @@ -334,16 +286,25 @@ static __global__ void conv2d_implicit_kernel(const float * __restrict__ input, } template -static void conv2d_implicit_cuda(const float * X_D, const T * K_D, float * Y_D, const conv_params P, cudaStream_t st) { - const int blocks = (P.TOTAL + CUDA_CONV2D_BLOCK_SIZE - 1) / CUDA_CONV2D_BLOCK_SIZE; - conv2d_implicit_kernel<<>>(X_D, K_D, Y_D, P); +static void conv2d_implicit_cuda(const float * X_D, const T * K_D, float * Y_D, const param_t &P, cudaStream_t st) { + // const int blocks = (P.TOTAL + CUDA_CONV2D_BLOCK_SIZE - 1) / CUDA_CONV2D_BLOCK_SIZE; + int blockx = ((P.Oh * P.Ow + 127) / 128); // blockx number + int blocky = (P.k + 127) / 128; // blocky number + int blockz = P.n; // blockz number + int threadx = CUDA_CONV2D_IMPLICT_BLOCK_SIZE; // threadx number per block + int thready = 1; // thready number per block + int threadz = 1; // threadz number per block + dim3 thblock(threadx, thready, threadz); + dim3 grid(blockx, blocky, blockz); + int smem_size = 24 * 1024; + conv2d_implicit_kernel<<>>(X_D, K_D, Y_D, P); } -static void conv2d_implicit_cuda_f16(const float * X_D, const half * K_D, float * Y_D, const conv_params P, cudaStream_t st) { +static void conv2d_implicit_cuda_f16(const float * X_D, const half * K_D, float * Y_D, const param_t &P, cudaStream_t st) { conv2d_implicit_cuda(X_D, K_D, Y_D, P, st); } -static void conv2d_implicit_cuda_f32(const float * X_D, const float * K_D, float * Y_D, const conv_params P, cudaStream_t st) { +static void conv2d_implicit_cuda_f32(const float * X_D, const float * K_D, float * Y_D, const param_t &P, cudaStream_t st) { conv2d_implicit_cuda(X_D, K_D, Y_D, P, st); } @@ -384,7 +345,8 @@ void ggml_cuda_op_conv2d_implicit(ggml_backend_cuda_context & ctx, ggml_tensor * const int B = input->ne[3]; // n_batches const int64_t total = B * OC * OH * OW; - conv_params params = { IW, IH, OW, OH, KW, KH, ST_X, ST_Y, PD_X, PD_Y, DL_X, DL_Y, IC, OC, B, total }; + // param_t params = { IW, IH, OW, OH, KW, KH, ST_X, ST_Y, PD_X, PD_Y, DL_X, DL_Y, IC, OC, B, total }; + param_t params = { B, IC, IH, IW, OC, KH, KW, ST_X, ST_Y, PD_X, PD_Y, DL_X, DL_Y, OH, OW }; if (kernel->type == GGML_TYPE_F16) { conv2d_implicit_cuda_f16(X_D, (half *) K_D, Y_D, params, st); diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index e06f95f081..0b799fbaf1 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -13,6 +13,7 @@ #include "ggml-cuda/concat.cuh" #include "ggml-cuda/conv-transpose-1d.cuh" #include "ggml-cuda/conv2d.cuh" +#include "ggml-cuda/conv2d-implicit.cuh" #include "ggml-cuda/conv2d-dw.cuh" #include "ggml-cuda/conv2d-transpose.cuh" #include "ggml-cuda/convert.cuh" @@ -2455,6 +2456,9 @@ static bool ggml_cuda_compute_forward(ggml_backend_cuda_context & ctx, struct gg case GGML_OP_CONV_2D: ggml_cuda_op_conv2d(ctx, dst); break; + case GGML_OP_CONV_2D_IMPLICIT: + ggml_cuda_op_conv2d_implicit(ctx, dst); + break; case GGML_OP_CONV_2D_DW: ggml_cuda_op_conv2d_dw(ctx, dst); break; @@ -3560,6 +3564,7 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g } case GGML_OP_IM2COL: case GGML_OP_CONV_2D: + case GGML_OP_CONV_2D_IMPLICIT: case GGML_OP_CONV_2D_DW: case GGML_OP_CONV_TRANSPOSE_2D: case GGML_OP_POOL_2D: diff --git a/ggml/src/ggml.c b/ggml/src/ggml.c index 4e0fd672bd..69003dfc5c 100644 --- a/ggml/src/ggml.c +++ b/ggml/src/ggml.c @@ -1018,7 +1018,7 @@ static const char * GGML_OP_NAME[GGML_OP_COUNT] = { "GLU", }; -static_assert(GGML_OP_COUNT == 89, "GGML_OP_COUNT != 89"); +static_assert(GGML_OP_COUNT == 90, "GGML_OP_COUNT != 90"); static const char * GGML_OP_SYMBOL[GGML_OP_COUNT] = { "none", @@ -1121,7 +1121,7 @@ static const char * GGML_OP_SYMBOL[GGML_OP_COUNT] = { "glu(x)", }; -static_assert(GGML_OP_COUNT == 89, "GGML_OP_COUNT != 89"); +static_assert(GGML_OP_COUNT == 90, "GGML_OP_COUNT != 90"); static_assert(GGML_OP_POOL_COUNT == 2, "GGML_OP_POOL_COUNT != 2");