Repository: ggml-org/llama.cpp
Pull request: CUDA: use FP32 arithmetic for conv2d

Fixes https://github.com/ggml-org/llama.cpp/issues/15680 .

The kernel can simply be made to always use FP32 arithmetic, this also has the advantage of not gimping performance on Pascal. Performance differences (on an RTX 4090) are negligible:

| Backend   | GGML op   | Op parameters                                                                                                                             |   TFLOPS ef476916b |   TFLOPS 2721aaf89 |   Speedup |
|:----------|:----------|:------------------------------------------------------------------------------------------------------------------------------------------|-------------------:|-------------------:|----------:|
| CUDA0     | CONV_2D   | ne_input=[16,16,128,8],ne_kernel=[3,3,128,512],type_kernel=f16,stride0=1,stride1=1,padding0=0,padding1=0,dilation0=1,dilation1=1,cwhn=0   |               1.93 |               1.94 |      1.00 |
| CUDA0     | CONV_2D   | ne_input=[16,16,128,8],ne_kernel=[3,3,128,512],type_kernel=f32,stride0=1,stride1=1,padding0=0,padding1=0,dilation0=1,dilation1=1,cwhn=0   |               2.06 |               2.06 |      1.00 |
| CUDA0     | CONV_2D   | ne_input=[19,19,256,16],ne_kernel=[4,4,256,4096],type_kernel=f16,stride0=1,stride1=1,padding0=0,padding1=0,dilation0=1,dilation1=1,cwhn=0 |               2.26 |               2.33 |      1.03 |
| CUDA0     | CONV_2D   | ne_input=[19,19,256,16],ne_kernel=[4,4,256,4096],type_kernel=f32,stride0=1,stride1=1,padding0=0,padding1=0,dilation0=1,dilation1=1,cwhn=0 |               2.36 |               2.37 |      1.00 |
| CUDA0     | CONV_2D   | ne_input=[19,19,4,16],ne_kernel=[2,2,4,4],type_kernel=f16,stride0=1,stride1=1,padding0=0,padding1=0,dilation0=1,dilation1=1,cwhn=0        |               0.16 |               0.16 |      0.99 |
| CUDA0     | CONV_2D   | ne_input=[19,19,4,16],ne_kernel=[2,2,4,4],type_kernel=f32,stride0=1,stride1=1,padding0=0,padding1=0,dilation0=1,dilation1=1,cwhn=0        |               0.16 |               0.16 |      1.00 |
| CUDA0     | CONV_2D   | ne_input=[19,19,8,16],ne_kernel=[4,4,8,128],type_kernel=f16,stride0=1,stride1=1,padding0=0,padding1=0,dilation0=1,dilation1=1,cwhn=0      |               1.87 |               1.93 |      1.03 |
| CUDA0     | CONV_2D   | ne_input=[19,19,8,16],ne_kernel=[4,4,8,128],type_kernel=f32,stride0=1,stride1=1,padding0=0,padding1=0,dilation0=1,dilation1=1,cwhn=0      |               1.94 |               1.94 |      1.00 |
| CUDA0     | CONV_2D   | ne_input=[19,19,8,16],ne_kernel=[4,4,8,130],type_kernel=f16,stride0=1,stride1=1,padding0=0,padding1=0,dilation0=1,dilation1=1,cwhn=0      |               1.81 |               1.86 |      1.03 |
| CUDA0     | CONV_2D   | ne_input=[19,19,8,16],ne_kernel=[4,4,8,130],type_kernel=f32,stride0=1,stride1=1,padding0=0,padding1=0,dilation0=1,dilation1=1,cwhn=0      |               1.88 |               1.88 |      1.00 |
| CUDA0     | CONV_2D   | ne_input=[224,224,1,1],ne_kernel=[2,2,1,8],type_kernel=f16,stride0=1,stride1=1,padding0=0,padding1=0,dilation0=1,dilation1=1,cwhn=0       |               0.26 |               0.26 |      1.00 |
| CUDA0     | CONV_2D   | ne_input=[224,224,1,1],ne_kernel=[2,2,1,8],type_kernel=f32,stride0=1,stride1=1,padding0=0,padding1=0,dilation0=1,dilation1=1,cwhn=0       |               0.27 |               0.27 |      1.00 |
| CUDA0     | CONV_2D   | ne_input=[224,224,1,8],ne_kernel=[2,2,1,8],type_kernel=f16,stride0=1,stride1=1,padding0=0,padding1=0,dilation0=1,dilation1=1,cwhn=0       |               0.31 |               0.32 |      1.00 |
| CUDA0     | CONV_2D   | ne_input=[224,224,1,8],ne_kernel=[2,2,1,8],type_kernel=f32,stride0=1,stride1=1,padding0=0,padding1=0,dilation0=1,dilation1=1,cwhn=0       |               0.32 |               0.32 |      1.00 |
| CUDA0     | CONV_2D   | ne_input=[224,224,3,1],ne_kernel=[3,3,3,8],type_kernel=f16,stride0=1,stride1=1,padding0=0,padding1=0,dilation0=1,dilation1=1,cwhn=0       |               1.05 |               1.05 |      1.01 |
| CUDA0     | CONV_2D   | ne_input=[224,224,3,1],ne_kernel=[3,3,3,8],type_kernel=f32,stride0=1,stride1=1,padding0=0,padding1=0,dilation0=1,dilation1=1,cwhn=0       |               1.09 |               1.09 |      1.00 |
| CUDA0     | CONV_2D   | ne_input=[58,58,32,1],ne_kernel=[3,3,32,64],type_kernel=f16,stride0=1,stride1=1,padding0=0,padding1=0,dilation0=1,dilation1=1,cwhn=0      |               1.49 |               1.49 |      1.00 |
| CUDA0     | CONV_2D   | ne_input=[58,58,32,1],ne_kernel=[3,3,32,64],type_kernel=f32,stride0=1,stride1=1,padding0=0,padding1=0,dilation0=1,dilation1=1,cwhn=0      |               1.57 |               1.57 |      1.00 |
| CUDA0     | CONV_2D   | ne_input=[58,58,32,8],ne_kernel=[3,3,32,64],type_kernel=f16,stride0=1,stride1=1,padding0=0,padding1=0,dilation0=1,dilation1=1,cwhn=0      |               1.89 |               1.90 |      1.01 |
| CUDA0     | CONV_2D   | ne_input=[58,58,32,8],ne_kernel=[3,3,32,64],type_kernel=f32,stride0=1,stride1=1,padding0=0,padding1=0,dilation0=1,dilation1=1,cwhn=0      |               2.00 |               2.00 |      1.00 |


Changed files:
filename: ggml/src/ggml-cuda/conv2d.cu
patch: @@ -82,7 +82,7 @@ static __global__ void conv2d_kernel(const float * __restrict__ input,
     int64_t n, c_out, out_y, out_x;
     Layout::unpack_indices(global_idx, P, n, c_out, out_y, out_x);
 
-    T acc = 0;
+    float acc = 0.0f;
 
     for (int64_t c_in = 0; c_in < P.IC; ++c_in) {
         kernel_bounds bounds = calculate_kernel_bounds(out_x, out_y, P);
@@ -93,21 +93,15 @@ static __global__ void conv2d_kernel(const float * __restrict__ input,
             for (int64_t kx = bounds.x_min; kx < bounds.x_max; ++kx) {
                 const int64_t in_x = calculate_input_coord(out_x, kx, P.ST_X, P.DL_X, P.PD_X);
 
-                T input_val;
-                if (std::is_same<T, half>::value) {
-                    input_val = __float2half(input[Layout::input_index(n, c_in, in_y, in_x, P)]);
-                } else {
-                    input_val = input[Layout::input_index(n, c_in, in_y, in_x, P)];
-                }
-
-                T kernel_val = kernel[Layout::kernel_index(c_out, c_in, ky, kx, P)];
+                const float input_val = input[Layout::input_index(n, c_in, in_y, in_x, P)];
+                const float kernel_val = kernel[Layout::kernel_index(c_out, c_in, ky, kx, P)];
                 acc += (input_val * kernel_val);
             }
         }
     }
 
     // [N, OC, OH, OW]
-    output[Layout::output_index(n, c_out, out_y, out_x, P)] = (float) acc;
+    output[Layout::output_index(n, c_out, out_y, out_x, P)] = acc;
 }
 
 template <typename T>