char-fork/llama.cpp

git clone https://git.t4t.associates/char-fork/llama.cpp

Andreas KieslingerAvoid PDL race conditions by disabling __restrict__ when PDL is used (#24030)9e58d4d69

char/aliceai-t5-moe
brancheschar/aliceai-t5-moeall refs →
1.2 KiB43 linesraw
1#include "common.cuh"
2
3// Row reduction kernel template - compute sum (norm=false) or mean (norm=true)
4template <bool norm>
5static __global__ void reduce_rows_f32(const float * x_ptr, float * dst_ptr, const int ncols) {
6    const float * GGML_CUDA_RESTRICT x   = x_ptr;
7    float       * GGML_CUDA_RESTRICT dst = dst_ptr;
8    const int row = blockIdx.x;
9    const int col = threadIdx.x;
10
11    float     sum        = 0.0f;
12    const int num_unroll = 8;
13    float     temp[num_unroll];
14    float     sum_temp[num_unroll] = { 0.0f };
15
16    ggml_cuda_pdl_sync();
17    for (int i = col; i < ncols;) {
18        for (int j = 0; j < num_unroll; ++j) {
19            if (i < ncols) {
20                temp[j] = x[row * ncols + i];
21            } else {
22                temp[j] = 0;
23            }
24            i += blockDim.x;
25        }
26        for (int j = 0; j < num_unroll; ++j) {
27            sum_temp[j] += temp[j];
28        }
29    }
30    for (int j = 0; j < num_unroll; ++j) {
31        sum += sum_temp[j];
32    }
33
34    // sum up partial sums
35    __shared__ float shared_vals[32];
36    sum = block_reduce<block_reduce_method::SUM>(sum, shared_vals);
37
38    if (col != 0) {
39        return;
40    }
41
42    dst[row] = norm ? sum / ncols : sum;
43}