Test CUDA conv2D type conversion fix

.github/workflows/menlo-build.yml

@@ -25,6 +25,9 @@ on:
       ]
   workflow_dispatch:
 
+  pull_request:
+    branches: ["dev"]
+
 env:
   VULKAN_VERSION: 1.3.261.1
 
@@ -56,7 +59,6 @@ jobs:
 
   build-and-test:
     runs-on: ${{ matrix.runs-on }}
-    needs: [create-draft-release]
     timeout-minutes: 270
     strategy:
       fail-fast: false
@@ -285,7 +287,7 @@ jobs:
         uses: actions/checkout@v3
         with:
           submodules: recursive
-      
+
       - name: Replace our Makefile
         run: |
           cat menlo/Makefile | tee Makefile
@@ -635,4 +637,4 @@ jobs:
           upload_url: ${{ needs.create-draft-release.outputs.upload_url }}
           asset_path: /tmp/cudart-llama-bin-win-cu11.7-x64.tar.gz
           asset_name: cudart-llama-bin-win-cu11.7-x64.tar.gz
-          asset_content_type: application/gzip
+          asset_content_type: application/gzip

ggml/src/ggml-cuda/conv2d.cu

@@ -1,4 +1,5 @@
 #include "conv2d.cuh"
+#include "convert.cuh"
 
 struct conv_params {
     const int64_t IW, IH;
@@ -94,8 +95,8 @@ static __global__ void conv2d_kernel(const float * __restrict__ input,
                 const int64_t in_x = calculate_input_coord(out_x, kx, P.ST_X, P.DL_X, P.PD_X);
 
                 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);
+                const T kernel_val = kernel[Layout::kernel_index(c_out, c_in, ky, kx, P)];
+                acc += (input_val * ggml_cuda_cast<float, T>(kernel_val));
             }
         }
     }