cuda: stage the IQ2_XXS codebook in shared memory on SM75 (experiment)

Nsight (via the CMP unlock) on mul_mat_vec_q<IQ2_XXS,ncols=1>, SM75:
  lg_throttle      19.81/issue   <- dominant, 43x the next stall
  wait              1.43
  long_scoreboard   0.46
  mio_throttle      0.06         <- shared/MIO pipe idle
  L1 hit 98.04%, DRAM 22.99%, SM 28.54%, occupancy 92.82%

Not bandwidth, not cache, not occupancy: the global load/store issue queue is
full. A vec_dot issues ~16 loads (4x16-bit from get_int_b2, 4x64-bit grid, 8x
32-bit activations); a load costs an LG issue slot even when it hits L1.

Move the four codebook lookups to LDS, which uses the MIO pipe that is sitting
idle. 2 KiB of shared per block, loaded once before the k-loop.

ggml/src/ggml-cuda/mmvq.cu

@@ -743,12 +743,47 @@ static __global__ void mul_mat_vec_q(
     const block_q8_1 * y = ((const block_q8_1 *) vy) + sample_y*stride_sample_y + channel_y*stride_channel_y;
     const int kbx_offset = sample_x*stride_sample_x + channel_x*stride_channel_x + row0*stride_row_x;
 
+#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ == GGML_CUDA_CC_TURING
+    // Stage the IQ2_XXS codebook in shared memory once per block. The kernel is
+    // not cache- or bandwidth-limited (L1 98%, DRAM 23%) but stalls on
+    // lg_throttle 19.81/issue: the global load/store issue queue is saturated
+    // while the shared/MIO pipe is idle at 0.06. Only 2 KiB, and it takes four
+    // LDG per vec_dot call off the contended pipe.
+    __shared__ uint64_t s_iq2xxs_grid[type == GGML_TYPE_IQ2_XXS ? 256 : 1];
+    if constexpr (type == GGML_TYPE_IQ2_XXS) {
+        for (int i = tid; i < 256; i += nwarps*warp_size) {
+            s_iq2xxs_grid[i] = iq2xxs_grid[i];
+        }
+        __syncthreads();
+    }
+#endif
+
     for (int kbx = tid / (qi/vdr); kbx < blocks_per_row_x; kbx += blocks_per_iter) {
         const int kby = kbx * (qk/QK8_1); // y block index that aligns with kbx
 
         // x block quant index when casting the quants to int
         const int kqs = vdr * (tid % (qi/vdr));
 
+#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ == GGML_CUDA_CC_TURING
+        if constexpr (type == GGML_TYPE_IQ2_XXS) {
+#pragma unroll
+            for (int j = 0; j < ncols_dst; ++j) {
+#pragma unroll
+                for (int i = 0; i < rows_per_cuda_block; ++i) {
+                    tmp[j][i] += vec_dot_iq2_xxs_q8_1_sgrid(
+                        vx, &y[j*stride_col_y + kby], kbx_offset + i*stride_row_x + kbx, kqs, s_iq2xxs_grid);
+                    if constexpr (has_fusion) {
+                        if (use_gate) {
+                            tmp_gate[j][i] += vec_dot_iq2_xxs_q8_1_sgrid(
+                                vgate, &y[j*stride_col_y + kby], kbx_offset + i*stride_row_x + kbx, kqs, s_iq2xxs_grid);
+                        }
+                    }
+                }
+            }
+            continue;
+        }
+#endif
+
 #if defined(__CUDA_ARCH__) && __CUDA_ARCH__ == GGML_CUDA_CC_TURING && GGML_SM75_IQ4_REUSE
         if constexpr (type == GGML_TYPE_IQ4_XS && ncols_dst >= 2 && ncols_dst <= 4) {
 #pragma unroll

ggml/src/ggml-cuda/vecdotq.cuh

@@ -1046,6 +1046,48 @@ static __device__ __forceinline__ float vec_dot_q6_K_q8_1(
 #define VDR_IQ2_XXS_Q8_1_MMVQ 2
 #define VDR_IQ2_XXS_Q8_1_MMQ  2
 
+#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ == GGML_CUDA_CC_TURING
+// Identical maths to vec_dot_iq2_xxs_q8_1, but the codebook is read from shared
+// memory. Nsight on SM75 shows this kernel stalling almost entirely on
+// lg_throttle (19.81 per issue, 43x the next reason) -- the load/store unit's
+// instruction queue is full, while the MIO/shared pipe is idle (mio_throttle
+// 0.06). L1 hit rate is already 98%, so this is not about cache misses: a global
+// load occupies an LG issue slot whether or not it hits. Turning the four grid
+// lookups into LDS moves them onto the pipe that has capacity.
+static __device__ __forceinline__ float vec_dot_iq2_xxs_q8_1_sgrid(
+    const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs,
+    const uint64_t * __restrict__ s_grid) {
+
+    const block_iq2_xxs * bq2 = (const block_iq2_xxs *) vbq + kbx;
+
+    const int q2 = get_int_b2(bq2->qs, iqs);
+    const uint8_t * aux8 = (const uint8_t *) &q2;
+    const uint32_t aux32 = get_int_b2(bq2->qs, iqs + 1);
+
+    int sumi = 0;
+#pragma unroll
+    for (int k0 = 0; k0 < 8; k0 += 2) {
+        const uint2    grid_pos = ((const uint2 *) s_grid)[aux8[k0/2]];
+        const uint32_t signs    = unpack_ksigns(aux32 >> (7 * k0 / 2));
+
+        const int signs0 = __vcmpne4(signs & 0x08040201, 0);
+        const int grid0  = __vsub4(grid_pos.x ^ signs0, signs0);
+        const int u0     = get_int_b4(bq8_1[iqs/2].qs, k0 + 0);
+        sumi = ggml_cuda_dp4a(grid0, u0, sumi);
+
+        const int signs1 = __vcmpne4(signs & 0x80402010, 0);
+        const int grid1  = __vsub4(grid_pos.y ^ signs1, signs1);
+        const int u1     = get_int_b4(bq8_1[iqs/2].qs, k0 + 1);
+        sumi = ggml_cuda_dp4a(grid1, u1, sumi);
+    }
+
+    const int ls = aux32 >> 27 | 1; // (scale * 2 + 1)
+    sumi = sumi * ls / 8;           // (sumi * scale + sumi / 2) / 4
+    const float d = __half2float(bq2->d) * __low2float(bq8_1[iqs/2].ds);
+    return d * sumi;
+}
+#endif
+
 static __device__ __forceinline__ float vec_dot_iq2_xxs_q8_1(
     const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs) {