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

Code before:
    *(float *)(dst_ptr + i0 * nb0) = value;
}

void ggml_cuda_op_pad_reflect_1d(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
    const ggml_tensor * src0 = dst->src[0];
    cudaStream_t stream = ctx.stream();

    GGML_ASSERT(src0->type == GGML_TYPE_F32);
    GGML_ASSERT(dst->type == GGML_TYPE_F32);

    const int32_t * opts = (const int32_t *) dst->op_params;
    const int p0 = opts[0];
    const int p1 = opts[1];

    const int64_t ne00 = src0->ne[0];
    const int64_t ne01 = src0->ne[1];
    const uint3   ne01_packed = init_fastdiv_values(ne01);
    const int64_t ne02 = src0->ne[2];
    const int64_t ne03 = src0->ne[3];

    const int64_t ne0 = dst->ne[0];

    // sanity: padded length matches
    GGML_ASSERT(ne0 == ne00 + p0 + p1);

    constexpr int64_t bx =  CUDA_PAD_REFLECT_1D_BLOCK_SIZE; // threads per block (x)
    const int64_t tiles0 = (ne0 + bx - 1) / bx; // number of tiles along i0
    // grid.x covers i1 and all tiles of i0: [ne01 * tiles0]
    // grid.y covers i2: [ne02]
    // grid.z covers i3: [ne03]
    const dim3 grid_dims((unsigned)(ne01 * tiles0), (unsigned)ne02, (unsigned)ne03);
    const dim3 block_dims((unsigned)bx, 1, 1);

    pad_reflect_1d_kernel_f32<<<grid_dims, block_dims, 0, stream>>>(
        src0->data, dst->data,
        ne0, ne00, ne01, ne01_packed, ne02, ne03,
        src0->nb[0], src0->nb[1], src0->nb[2], src0->nb[3],
        dst->nb[0], dst->nb[1], dst->nb[2], dst->nb[3],
        p0, p1
    );
}


Review comment (suggestion):
```suggestion
    constexpr int64_t bx = CUDA_PAD_REFLECT_1D_BLOCK_SIZE; // threads per block (x)
```

Code after:
    const float value               = *(const float *) (src0_ptr + src_idx * nb00);
    *(float *) (dst_ptr + i0 * nb0) = value;
}

void ggml_cuda_op_pad_reflect_1d(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
    const ggml_tensor * src0   = dst->src[0];
    cudaStream_t        stream = ctx.stream();

    GGML_ASSERT(src0->type == GGML_TYPE_F32);
    GGML_ASSERT(dst->type == GGML_TYPE_F32);

    const int32_t * opts = (const int32_t *) dst->op_params;
    const int       p0   = opts[0];
    const int       p1   = opts[1];

    const int64_t ne00        = src0->ne[0];
    const int64_t ne01        = src0->ne[1];
    const uint3   ne01_packed = init_fastdiv_values(ne01);
    const int64_t ne02        = src0->ne[2];
    const int64_t ne03        = src0->ne[3];

    const int64_t ne0 = dst->ne[0];

    // sanity: padded length matches
    GGML_ASSERT(ne0 == ne00 + p0 + p1);

    constexpr int64_t bx     = CUDA_PAD_REFLECT_1D_BLOCK_SIZE;  // threads per block (x)
    const int64_t     tiles0 = (ne0 + bx - 1) / bx;             // number of tiles along i0
    // grid.x covers i1 and all tiles of i0: [ne01 * tiles0]
    // grid.y covers i2: [ne02]
    // grid.z covers i3: [ne03]
    const dim3        grid_dims((unsigned) (ne01 * tiles0), (unsigned) ne02, (unsigned) ne03);
    const dim3        block_dims((unsigned) bx, 1, 1);

    pad_reflect_1d_kernel_f32<<<grid_dims, block_dims, 0, stream>>>(
        src0->data, dst->data, ne0, ne00, ne01_packed, ne02, ne03, src0->nb[0], src0->nb[1], src0->nb[2], src0->nb[3],
        dst->nb[0], dst->nb[1], dst->nb[2], dst->nb[3], p0, p1);
}