diff --git a/ggml/src/ggml-cuda/mma.cuh b/ggml/src/ggml-cuda/mma.cuh
index 6af2b6a14..3583ba5e1 100644
--- a/ggml/src/ggml-cuda/mma.cuh
+++ b/ggml/src/ggml-cuda/mma.cuh
@@ -873,14 +873,16 @@ namespace ggml_cuda_mma {
     }
 
     // Load from tile element (i0, j0), swz tells if the tile is stored swizzled.
-    template <bool swz, typename T, data_layout dl>
+    template <bool swz, int I, int J, typename T, data_layout dl>
     static __device__ __forceinline__ void load_ldmatrix(
-            tile<16, 8, T, dl> & t, const T * __restrict__ tile_base, const int i0, const int j0, const int stride) {
+            tile<I, J, T, dl> & t, const T * __restrict__ tile_base, const int i0, const int j0, const int stride) {
         if constexpr (!swz) {
             load_ldmatrix(t, tile_base + i0*stride + j0, stride);
             return;
         }
 #if defined(TURING_MMA_AVAILABLE)
+        static_assert(I == 16, "bad tile width");
+        static_assert(J ==  8, "bad tile height");
         const int i = i0 + threadIdx.x % t.I;
         const int j = j0 + (threadIdx.x / t.I) * (t.J / 2);
         int * xi = (int *) t.x;