Skip to content

Commit 0a524f2

Browse files
AgainstEntropyam17an
AgainstEntropy
and
am17an
authored
CUDA & CPU: support F32 kernel type for CONV_TRANSPOSE_2D (#17094)
* Refactor CUDA 2D transpose implementation to support multiple kernel types and improve parameter handling - Introduced a `conv2d_transpose_params` struct for better parameter management. - Updated `conv2d_transpose_kernel` to be templated for different kernel types (float and half). - Modified `ggml_cuda_conv_2d_transpose_p0` to handle both F16 and F32 kernel types. - Enhanced test cases to validate functionality for both kernel types. * Refactor test cases for 2D convolution transpose to support dynamic kernel types - Updated `test_conv_transpose_2d` structure to improve parameter handling by reordering constructor arguments. - Enhanced test case generation to iterate over kernel types, allowing for flexible testing of different configurations. - Removed hardcoded kernel type instances in favor of a loop for better maintainability and scalability. * Refactor ggml_compute_forward_conv_transpose_2d to support both F16 and F32 tensor types. * Refactor conv2d transpose kernel to use a template for kernel type, enhancing flexibility for different data types. Update test cases to include both F16 and F32 tensor types for comprehensive coverage. * Update ggml/src/ggml-cuda/conv2d-transpose.cu Co-authored-by: Aman Gupta <amangupta052@gmail.com> * Update ggml/src/ggml-cpu/ggml-cpu.c Co-authored-by: Aman Gupta <amangupta052@gmail.com> * Refactor conv2d transpose implementation by removing the conv2d_transpose_params struct and dispatching with direct kernel launch. * Enhance cpu conv2d transpose implementation by introducing a templated kernel type for improved flexibility with F16 and F32 data types. --------- Co-authored-by: Aman Gupta <amangupta052@gmail.com>
1 parent c0159f9 commit 0a524f2

5 files changed

Lines changed: 123 additions & 54 deletions

File tree

‎ggml/src/ggml-cpu/ggml-cpu.c‎

Lines changed: 6 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -2871,8 +2871,12 @@ struct ggml_cplan ggml_graph_plan(
28712871
const int64_t ne11 = node->src[1]->ne[1]; // H
28722872
const int64_t ne12 = node->src[1]->ne[2]; // Channels In
28732873

2874-
cur += sizeof(ggml_fp16_t)*ne00*ne01*ne02*ne03;
2875-
cur += sizeof(ggml_fp16_t)*ne10*ne11*ne12;
2874+
GGML_ASSERT(node->src[0]->type == GGML_TYPE_F16 || node->src[0]->type == GGML_TYPE_F32);
2875+
GGML_ASSERT(node->src[1]->type == GGML_TYPE_F32);
2876+
2877+
cur += ggml_type_size(node->src[0]->type) * ne00 * ne01 * ne02 * ne03;
2878+
cur += ggml_type_size(node->src[0]->type) * ne10 * ne11 * ne12;
2879+
28762880
} break;
28772881
case GGML_OP_TOP_K:
28782882
{

‎ggml/src/ggml-cpu/ops.cpp‎

Lines changed: 50 additions & 19 deletions
Original file line numberDiff line numberDiff line change
@@ -6923,16 +6923,15 @@ void ggml_compute_forward_conv_3d(
69236923
ggml_compute_forward_conv_3d_impl(params, src0, src1, dst, src0->type);
69246924
}
69256925

6926-
// ggml_compute_forward_conv_transpose_2d
6927-
6928-
void ggml_compute_forward_conv_transpose_2d(
6929-
const ggml_compute_params * params,
6930-
ggml_tensor * dst) {
6926+
template <typename kernel_t>
6927+
static void ggml_compute_forward_conv_transpose_2d_impl(
6928+
const ggml_compute_params * params,
6929+
ggml_tensor * dst) {
69316930

69326931
const ggml_tensor * src0 = dst->src[0];
69336932
const ggml_tensor * src1 = dst->src[1];
69346933

6935-
GGML_ASSERT(src0->type == GGML_TYPE_F16);
6934+
GGML_ASSERT(src0->type == GGML_TYPE_F16 || src0->type == GGML_TYPE_F32);
69366935
GGML_ASSERT(src1->type == GGML_TYPE_F32);
69376936
GGML_ASSERT( dst->type == GGML_TYPE_F32);
69386937

@@ -6943,20 +6942,20 @@ void ggml_compute_forward_conv_transpose_2d(
69436942

69446943
const int nk = ne00*ne01*ne02*ne03;
69456944

6946-
GGML_ASSERT(nb00 == sizeof(ggml_fp16_t));
6945+
GGML_ASSERT(nb00 == ggml_type_size(src0->type));
69476946
GGML_ASSERT(nb10 == sizeof(float));
69486947

69496948
if (ith == 0) {
69506949
memset(params->wdata, 0, params->wsize);
69516950

69526951
// permute kernel data (src0) from (Kw x Kh x Cout x Cin) to (Cin x Kw x Kh x Cout)
69536952
{
6954-
ggml_fp16_t * const wdata = (ggml_fp16_t *) params->wdata + 0;
6953+
kernel_t * const wdata = (kernel_t *) params->wdata + 0;
69556954

69566955
for (int64_t i03 = 0; i03 < ne03; i03++) {
69576956
for (int64_t i02 = 0; i02 < ne02; i02++) {
6958-
const ggml_fp16_t * const src = (ggml_fp16_t *)((char *) src0->data + i03*nb03 + i02*nb02);
6959-
ggml_fp16_t * dst_data = wdata + i02*ne01*ne00*ne03;
6957+
const kernel_t * const src = (kernel_t *)((char *) src0->data + i03*nb03 + i02*nb02);
6958+
kernel_t * dst_data = wdata + i02*ne01*ne00*ne03;
69606959
for (int64_t i01 = 0; i01 < ne01; i01++) {
69616960
for (int64_t i00 = 0; i00 < ne00; i00++) {
69626961
dst_data[i01*ne00*ne03 + i00*ne03 + i03] = src[i01 * ne00 + i00];
@@ -6968,13 +6967,17 @@ void ggml_compute_forward_conv_transpose_2d(
69686967

69696968
// permute source data (src1) from (Sw x Sh x Cin) to (Cin x Sw x Sh)
69706969
{
6971-
ggml_fp16_t * const wdata = (ggml_fp16_t *) params->wdata + nk;
6970+
kernel_t * const wdata = (kernel_t *) params->wdata + nk;
69726971
for (int i12 = 0; i12 < ne12; i12++) {
69736972
for (int i11 = 0; i11 < ne11; i11++) {
69746973
const float * const src = (float *)((char *) src1->data + i12*nb12 + i11*nb11);
6975-
ggml_fp16_t * dst_data = wdata + i11*ne10*ne12;
6974+
kernel_t * dst_data = wdata + i11*ne10*ne12;
69766975
for (int i10 = 0; i10 < ne10; i10++) {
6977-
dst_data[i10*ne12 + i12] = GGML_CPU_FP32_TO_FP16(src[i10]);
6976+
if constexpr (std::is_same_v<kernel_t, ggml_fp16_t>) {
6977+
dst_data[i10*ne12 + i12] = GGML_CPU_FP32_TO_FP16(src[i10]);
6978+
} else {
6979+
dst_data[i10*ne12 + i12] = src[i10];
6980+
}
69786981
}
69796982
}
69806983
}
@@ -6996,21 +6999,27 @@ void ggml_compute_forward_conv_transpose_2d(
69966999
const int ip0 = dp*ith;
69977000
const int ip1 = MIN(ip0 + dp, np);
69987001

6999-
ggml_fp16_t * const wdata = (ggml_fp16_t *) params->wdata + 0;
7000-
ggml_fp16_t * const wdata_src = wdata + nk;
7002+
kernel_t * const wdata = (kernel_t *) params->wdata + 0;
7003+
kernel_t * const wdata_src = wdata + nk;
70017004

70027005
for (int i2 = ip0; i2 < ip1; i2++) { // Cout
70037006
float * dst_data = (float *)((char *) dst->data + i2*nb2);
7004-
ggml_fp16_t * wdata_kernel = wdata + i2*ne01*ne00*ne03;
7007+
kernel_t * wdata_kernel = wdata + i2*ne01*ne00*ne03;
70057008
for (int i11 = 0; i11 < ne11; i11++) {
70067009
for (int i10 = 0; i10 < ne10; i10++) {
70077010
const int i1n = i11*ne10*ne12 + i10*ne12;
70087011
for (int i01 = 0; i01 < ne01; i01++) {
70097012
for (int i00 = 0; i00 < ne00; i00++) {
70107013
float v = 0;
7011-
ggml_vec_dot_f16(ne03, &v, 0,
7012-
wdata_src + i1n, 0,
7013-
wdata_kernel + i01*ne00*ne03 + i00*ne03, 0, 1);
7014+
if constexpr (std::is_same_v<kernel_t, ggml_fp16_t>) {
7015+
ggml_vec_dot_f16(ne03, &v, 0,
7016+
wdata_src + i1n, 0,
7017+
wdata_kernel + i01*ne00*ne03 + i00*ne03, 0, 1);
7018+
} else {
7019+
ggml_vec_dot_f32(ne03, &v, 0,
7020+
wdata_src + i1n, 0,
7021+
wdata_kernel + i01*ne00*ne03 + i00*ne03, 0, 1);
7022+
}
70147023
dst_data[(i11*stride + i01)*ne0 + i10*stride + i00] += v;
70157024
}
70167025
}
@@ -7019,6 +7028,28 @@ void ggml_compute_forward_conv_transpose_2d(
70197028
}
70207029
}
70217030

7031+
void ggml_compute_forward_conv_transpose_2d(
7032+
const ggml_compute_params * params,
7033+
ggml_tensor * dst) {
7034+
7035+
const ggml_tensor * src0 = dst->src[0];
7036+
7037+
switch (src0->type) {
7038+
case GGML_TYPE_F16:
7039+
{
7040+
ggml_compute_forward_conv_transpose_2d_impl<ggml_fp16_t>(params, dst);
7041+
} break;
7042+
case GGML_TYPE_F32:
7043+
{
7044+
ggml_compute_forward_conv_transpose_2d_impl<float>(params, dst);
7045+
} break;
7046+
default:
7047+
{
7048+
GGML_ABORT("fatal error");
7049+
}
7050+
}
7051+
}
7052+
70227053
// ggml_compute_forward_conv_2d_dw
70237054

70247055
struct ggml_conv_2d_dw_params {

‎ggml/src/ggml-cuda/conv2d-transpose.cu‎

Lines changed: 45 additions & 21 deletions
Original file line numberDiff line numberDiff line change
@@ -1,12 +1,20 @@
1-
#include <algorithm>
2-
31
#include "conv2d-transpose.cuh"
4-
#include "ggml.h"
5-
6-
__global__ void conv2d_transpose_kernel(const float * __restrict__ input, const half * __restrict__ kernel,
7-
float * __restrict__ output, const int in_w, const int in_h, const int out_w,
8-
const int out_h, const int kernel_w, const int kernel_h, const int stride,
9-
const int c_in, const int c_out, const int batches) {
2+
#include "convert.cuh"
3+
4+
template <typename kernel_t>
5+
static __global__ void conv2d_transpose_kernel(const float * __restrict__ input,
6+
const kernel_t * __restrict__ kernel,
7+
float * __restrict__ output,
8+
const int in_w,
9+
const int in_h,
10+
const int out_w,
11+
const int out_h,
12+
const int kernel_w,
13+
const int kernel_h,
14+
const int stride,
15+
const int c_in,
16+
const int c_out,
17+
const int batches) {
1018
const int global_idx = blockIdx.x * blockDim.x + threadIdx.x;
1119

1220
const int total_elements = out_w * out_h * c_out * batches;
@@ -26,24 +34,32 @@ __global__ void conv2d_transpose_kernel(const float * __restrict__ input, const
2634
for (int c_in_idx = 0; c_in_idx < c_in; c_in_idx++) {
2735
for (int kh = 0; kh < kernel_h; ++kh) {
2836
int in_y = out_y_idx - kh;
29-
if (in_y < 0 || in_y % stride) continue;
37+
if (in_y < 0 || in_y % stride) {
38+
continue;
39+
}
3040
in_y /= stride;
31-
if (in_y >= in_h) continue;
41+
if (in_y >= in_h) {
42+
continue;
43+
}
3244

3345
for (int kw = 0; kw < kernel_w; ++kw) {
3446
int in_x = out_x_idx - kw;
35-
if (in_x < 0 || in_x % stride) continue;
47+
if (in_x < 0 || in_x % stride) {
48+
continue;
49+
}
3650
in_x /= stride;
37-
if (in_x >= in_w) continue;
51+
if (in_x >= in_w) {
52+
continue;
53+
}
3854

3955
const int input_idx = (in_w * in_h * c_in) * n_idx + (in_w * in_h) * c_in_idx + (in_w) *in_y + in_x;
4056
const int kernel_idx =
4157
(kernel_h * kernel_w * c_out) * c_in_idx + (kernel_h * kernel_w) * c_idx + (kernel_w) *kh + kw;
4258

43-
float input_val = input[input_idx];
44-
half kern_val = kernel[kernel_idx];
59+
float input_val = input[input_idx];
60+
kernel_t kern_val = kernel[kernel_idx];
4561

46-
accumulator += input_val * (float) kern_val;
62+
accumulator += input_val * ggml_cuda_cast<float>(kern_val);
4763
}
4864
}
4965
}
@@ -56,11 +72,12 @@ void ggml_cuda_conv_2d_transpose_p0(ggml_backend_cuda_context & ctx, ggml_tensor
5672
const ggml_tensor * kernel = dst->src[0];
5773
const ggml_tensor * input = dst->src[1];
5874

59-
GGML_ASSERT(kernel->type == GGML_TYPE_F16 && input->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F32);
75+
GGML_ASSERT(kernel->type == GGML_TYPE_F16 || kernel->type == GGML_TYPE_F32);
76+
GGML_ASSERT(input->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F32);
6077

6178
const float * input_data = (const float *) input->data;
6279
float * output_data = (float *) dst->data;
63-
const half * kernel_data = (const half *) kernel->data;
80+
const void * kernel_data = kernel->data;
6481

6582
const int input_w = input->ne[0];
6683
const int input_h = input->ne[1];
@@ -82,10 +99,17 @@ void ggml_cuda_conv_2d_transpose_p0(ggml_backend_cuda_context & ctx, ggml_tensor
8299
GGML_ASSERT(ggml_is_contiguous(kernel));
83100
GGML_ASSERT(ggml_is_contiguous(dst));
84101

85-
const int total = (output_w * output_h * channels_out * batches);
102+
const int total = output_w * output_h * channels_out * batches;
86103
const int blocks = (total + CUDA_CONV2D_TRANSPOSE_BLOCK_SIZE - 1) / CUDA_CONV2D_TRANSPOSE_BLOCK_SIZE;
87104

88-
conv2d_transpose_kernel<<<blocks, CUDA_CONV2D_TRANSPOSE_BLOCK_SIZE, 0, st>>>(
89-
input_data, kernel_data, output_data, input_w, input_h, output_w, output_h, kernel_w, kernel_h, stride,
90-
channels_in, channels_out, batches);
105+
if (kernel->type == GGML_TYPE_F16) {
106+
conv2d_transpose_kernel<half><<<blocks, CUDA_CONV2D_TRANSPOSE_BLOCK_SIZE, 0, st>>>(
107+
input_data, (const half *) kernel_data, output_data, input_w, input_h, output_w, output_h, kernel_w,
108+
kernel_h, stride, channels_in, channels_out, batches);
109+
110+
} else {
111+
conv2d_transpose_kernel<float><<<blocks, CUDA_CONV2D_TRANSPOSE_BLOCK_SIZE, 0, st>>>(
112+
input_data, (const float *) kernel_data, output_data, input_w, input_h, output_w, output_h, kernel_w,
113+
kernel_h, stride, channels_in, channels_out, batches);
114+
}
91115
}

‎ggml/src/ggml-cuda/conv2d-transpose.cuh‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,5 @@
11
#include "common.cuh"
22

33
#define CUDA_CONV2D_TRANSPOSE_BLOCK_SIZE 256
4+
45
void ggml_cuda_conv_2d_transpose_p0(ggml_backend_cuda_context & ctx, ggml_tensor * dst);

‎tests/test-backend-ops.cpp‎

Lines changed: 21 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -4823,28 +4823,33 @@ struct test_conv_transpose_1d : public test_case {
48234823

48244824
// GGML_OP_CONV_TRANSPOSE_2D
48254825
struct test_conv_transpose_2d : public test_case {
4826+
// Dimensions
48264827
const std::array<int64_t, 4> ne_input;
48274828
const std::array<int64_t, 4> ne_kernel;
48284829
const int stride;
4830+
// Types
4831+
const ggml_type kernel_type;
48294832

48304833
std::string vars() override {
4831-
return VARS_TO_STR3(ne_input, ne_kernel, stride);
4834+
return VARS_TO_STR4(kernel_type, ne_input, ne_kernel, stride);
48324835
}
48334836

48344837
double max_nmse_err() override {
48354838
return 5e-4; // The default 1e-7 is too small for Vulkan.
48364839
}
48374840

4838-
test_conv_transpose_2d(std::array<int64_t, 4> ne_input = {10, 10, 3, 1}, // [input_width, input_height, input_channels, 1]
4839-
std::array<int64_t, 4> ne_kernel = {3, 3, 3, 1}, // [kernel_width, kernel_height, input_channels, 1]
4840-
int stride = 1)
4841-
: ne_input(ne_input), ne_kernel(ne_kernel), stride(stride){}
4841+
test_conv_transpose_2d(
4842+
std::array<int64_t, 4> ne_input = {10, 10, 3, 1}, // [input_width, input_height, input_channels, 1]
4843+
std::array<int64_t, 4> ne_kernel = {3, 3, 3, 1}, // [kernel_width, kernel_height, input_channels, 1]
4844+
int stride = 1,
4845+
ggml_type kernel_type = GGML_TYPE_F16
4846+
) : ne_input(ne_input), ne_kernel(ne_kernel), stride(stride), kernel_type(kernel_type) {}
48424847

48434848
ggml_tensor * build_graph(ggml_context * ctx) override {
48444849
ggml_tensor * input = ggml_new_tensor(ctx, GGML_TYPE_F32, 4, ne_input.data());
48454850
ggml_set_name(input, "input");
48464851

4847-
ggml_tensor * kernel = ggml_new_tensor(ctx, GGML_TYPE_F16, 4, ne_kernel.data());
4852+
ggml_tensor * kernel = ggml_new_tensor(ctx, kernel_type, 4, ne_kernel.data());
48484853
ggml_set_name(kernel, "kernel");
48494854

48504855
ggml_tensor * out = ggml_conv_transpose_2d_p0(ctx, kernel, input, stride);
@@ -7704,9 +7709,11 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
77047709
test_cases.emplace_back(new test_conv_transpose_1d({3,2,1,1}, {3,1,2,1}, 1, 0, 1));
77057710
test_cases.emplace_back(new test_conv_transpose_1d({2,1,1,1}, {3,1,1,1}, 1, 0, 1));
77067711

7707-
test_cases.emplace_back(new test_conv_transpose_2d({3, 2, 3, 1}, {2, 2, 1, 3}, 1));
7708-
test_cases.emplace_back(new test_conv_transpose_2d({10, 10, 9, 1}, {3, 3, 1, 9}, 2));
7709-
test_cases.emplace_back(new test_conv_transpose_2d({129, 63, 35, 1}, {3, 3, 48, 35}, 1));
7712+
for (ggml_type kernel_type : {GGML_TYPE_F32, GGML_TYPE_F16}) {
7713+
test_cases.emplace_back(new test_conv_transpose_2d({3, 2, 3, 1}, {2, 2, 1, 3}, 1, kernel_type));
7714+
test_cases.emplace_back(new test_conv_transpose_2d({10, 10, 9, 1}, {3, 3, 1, 9}, 2, kernel_type));
7715+
test_cases.emplace_back(new test_conv_transpose_2d({129, 63, 35, 1}, {3, 3, 48, 35}, 1, kernel_type));
7716+
}
77107717

77117718
test_cases.emplace_back(new test_count_equal(GGML_TYPE_F32, {4, 500, 1, 1}));
77127719
test_cases.emplace_back(new test_count_equal(GGML_TYPE_F32, {4, 5000, 1, 1}));
@@ -8892,9 +8899,11 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_perf() {
88928899
test_cases.emplace_back(new test_conv_2d_dw({512, 512, 256, 1}, {3, 3, 1, 256}, 1, 1, 1, false));
88938900
test_cases.emplace_back(new test_conv_2d_dw({512, 512, 256, 1}, {3, 3, 1, 256}, 1, 1, 1, true));
88948901

8895-
test_cases.emplace_back(new test_conv_transpose_2d({256, 256, 256, 1}, {3, 3, 16, 256}, 1));
8896-
test_cases.emplace_back(new test_conv_transpose_2d({16, 16, 16, 1}, {3, 3, 8, 16}, 1));
8897-
test_cases.emplace_back(new test_conv_transpose_2d({10, 10, 9, 1}, {3, 3, 1, 9}, 2));
8902+
for (ggml_type kernel_type : {GGML_TYPE_F32, GGML_TYPE_F16}) {
8903+
test_cases.emplace_back(new test_conv_transpose_2d({256, 256, 256, 1}, {3, 3, 16, 256}, 1, kernel_type));
8904+
test_cases.emplace_back(new test_conv_transpose_2d({16, 16, 16, 1}, {3, 3, 8, 16}, 1, kernel_type));
8905+
test_cases.emplace_back(new test_conv_transpose_2d({10, 10, 9, 1}, {3, 3, 1, 9}, 2, kernel_type));
8906+
}
88988907

88998908
test_cases.emplace_back(new test_mean(GGML_TYPE_F32, {256, 256, 3, 1}));
89008909

0 commit comments

Comments
 (0)