Repository: ggerganov/llama.cpp
File: ggml/src/ggml-cuda/binbcast.cu (CUDA)

Code before:
                                           const uint3            ne0_fastdiv,
                                           const uint3            ne1_fastdiv,
                                           const uint3            ne2_fastdiv,
                                           const uint32_t         ne3,
                                           const uint3            ne012_fastdiv,
                                           const uint3            ne01_fastdiv,
                                           const uint3            ne10_fastdiv,
                                           const uint3            ne11_fastdiv,
                                           const uint3            ne12_fastdiv,
                                           const uint3            ne13_fastdiv,
                                           /*int s0, */ const int s1,
                                           const int              s2,
                                           const int              s3,
                                           /*int s00,*/ const int s01,
                                           const int              s02,
                                           const int              s03,
                                           /*int s10,*/ const int s11,
                                           const int              s12,
                                           const int              s13,
                                           src1_ptrs... src1s) {
    const int i = blockDim.x*blockIdx.x + threadIdx.x;

    const uint32_t i3 = fastdiv(i, ne012_fastdiv);
    const uint32_t i2 = fastmodulo(fastdiv(i, ne01_fastdiv), ne2_fastdiv);
    const uint32_t i1 = fastmodulo(fastdiv(i, ne0_fastdiv), ne1_fastdiv);
    const uint32_t i0 = fastmodulo(i, ne0_fastdiv);

    if (i0 >= ne0_fastdiv.z || i1 >= ne1_fastdiv.z || i2 >= ne2_fastdiv.z || i3 >= ne3) {
        return;
    }

    const int i11 = fastmodulo(i1, ne11_fastdiv);
    const int i12 = fastmodulo(i2, ne12_fastdiv);
    const int i13 = fastmodulo(i3, ne13_fastdiv);

    const size_t i_src0 =  i3*s03 +  i2*s02 +  i1*s01;
    const size_t i_src1 = i13*s13 + i12*s12 + i11*s11;
    const size_t i_dst  =  i3*s3  +  i2*s2  +  i1*s1;

    const src0_t * src0_row = src0 ? (src0 + i_src0) : nullptr;
    dst_t * dst_row = dst + i_dst;

    const int i10 = fastmodulo(i0, ne10_fastdiv);

    float result = src0_row ? (float) src0_row[i0] : 0.0f;
    if constexpr (sizeof...(src1_ptrs) > 0) {
        result = (..., (result = bin_op(result, (float)src1s[i_src1 + i10])));
    } else {
        result = bin_op(result, (float)src1[i_src1 + i10]);
    }


Review comment (suggestion):
```suggestion
    const uint32_t i3 = fastdiv(i, ne012_fastdiv);
    const uint32_t i2 = fastdiv(i - i3*ne012_fastdiv.z, ne01_fastdiv);
    const uint32_t i1 = fastdiv(i - i3*ne012_fastdiv.z - i2*ne01_fastdiv.z, ne0_fastdiv);
    const uint32_t i0 = i - i3*ne012_fastdiv.z - i2*ne01_fastdiv.z - i1*ne0_fastdiv.z;
```

The way the original code was written is inconsistent with how it's usually done in other parts of the code. The above code would be the usual pattern but with `fastdiv`.

Names with prefixes like `ne01` are confusing because that usually means index 1 of `src0`. Please either use the name without the `_fastdiv` for fastdiv values that are directly equivalent to tensor dimensions or a `prod_` prefix if they're equivalent to a product of tensor dimensions.

Code after:
                                           const uint3            ne0,
                                           const uint3            ne1,
                                           const uint3            ne2,
                                           const uint32_t         ne3,
                                           const uint3            prod_012,
                                           const uint3            prod_01,
                                           const uint3            ne10,
                                           const uint3            ne11,
                                           const uint3            ne12,
                                           const uint3            ne13,
                                           /*int s0, */ const int s1,
                                           const int              s2,
                                           const int              s3,
                                           /*int s00,*/ const int s01,
                                           const int              s02,
                                           const int              s03,
                                           /*int s10,*/ const int s11,
                                           const int              s12,
                                           const int              s13,
                                           src1_ptrs... src1s) {
    const int i = blockDim.x*blockIdx.x + threadIdx.x;

    const uint32_t i3 = fastdiv(i, prod_012);
    const uint32_t i2 = fastdiv(i - i3 * prod_012.z, prod_01);
    const uint32_t i1 = fastdiv(i - i3 * prod_012.z - i2 * prod_01.z, ne0);
    const uint32_t i0 = i - i3 * prod_012.z - i2 * prod_01.z - i1 * ne0.z;

    if (i0 >= ne0.z || i1 >= ne1.z || i2 >= ne2.z || i3 >= ne3) {
        return;
    }

    const int i11 = fastmodulo(i1, ne11);
    const int i12 = fastmodulo(i2, ne12);
    const int i13 = fastmodulo(i3, ne13);

    const size_t i_src0 =  i3*s03 +  i2*s02 +  i1*s01;
    const size_t i_src1 = i13*s13 + i12*s12 + i11*s11;
    const size_t i_dst  =  i3*s3  +  i2*s2  +  i1*s1;

    const src0_t * src0_row = src0 ? (src0 + i_src0) : nullptr;
    dst_t * dst_row = dst + i_dst;

    const int i10 = fastmodulo(i0, ne10);

    float result = src0_row ? (float) src0_row[i0] : 0.0f;
    if constexpr (sizeof...(src1_ptrs) > 0) {
        result = (..., (result = bin_op(result, (float)src1s[i_src1 + i10])));
    } else {
        result = bin_op(result, (float)src1[i_src1 + i10]);
    }