ronantakizawa/codereview-bench huggingface dataset:
Repository: ggerganov/llama.cpp
File: tools/server/webui/src/lib/components/app/chat/ChatSettings/ChatSettingsImportExportTab.svelte (Svelte)

Code before:
			const selectedData = $state
				.snapshot(fullImportData)
				.filter((item) => selectedIds.has(item.conv.id));

			await DatabaseStore.importConversations(selectedData);

			await chatStore.loadConversations();

			importedConversations = selectedConversations;
			showImportSummary = true;
			showExportSummary = false;
			showImportDialog = false;
		} catch (err) {
			console.error('Import failed:', err);
			alert('Failed to import conversations. Please check the file format.');
		}
	}

	async function handleDeleteAllClick() {
		try {
			const allConversations = await DatabaseStore.getAllConversations();
			if (allConversations.length === 0) {
				alert('No conversations to delete');
				return;
			}
			showDeleteDialog = true;
		} catch (err) {
			console.error('Failed to load conversations for deletion:', err);
			alert('Failed to load conversations');
		}
	}

	async function handleDeleteAllConfirm() {
		try {
			const allConversations = await DatabaseStore.getAllConversations();
			let deleted = 0;

			// Delete each conversation and its messages
			for (const conv of allConversations) {
				await DatabaseStore.deleteConversation(conv.id);
				deleted++;
			}

			// Update chat store to reflect the changes
			await chatStore.loadConversations();
			chatStore.clearActiveConversation();

			deletedCount = deleted;
			showDeleteSummary = true;
			showDeleteDialog = false;
			showExportSummary = false;

Review comment (suggestion):
```suggestion
			const allConversations = await DatabaseStore.getAllConversations();

			if (allConversations.length === 0) {
				alert('No conversations to delete');
				return;
			}

			showDeleteDialog = true;
```

Code after:

			await conversationsStore.importConversationsData(selectedData);

			importedConversations = selectedConversations;
			showImportSummary = true;
			showExportSummary = false;
			showImportDialog = false;
		} catch (err) {
			console.error('Import failed:', err);
			alert('Failed to import conversations. Please check the file format.');
		}
	}

	async function handleDeleteAllClick() {
		try {
			const allConversations = conversations();

			if (allConversations.length === 0) {
				toast.info('No conversations to delete');
				return;
			}

			showDeleteDialog = true;
		} catch (err) {
			console.error('Failed to load conversations for deletion:', err);
			toast.error('Failed to load conversations');
		}
	}

	async function handleDeleteAllConfirm() {
		try {
			await conversationsStore.deleteAll();

			showDeleteDialog = false;
		} catch (err) {
			console.error('Failed to delete conversations:', err);
		}
	}

	function handleDeleteAllCancel() {
		showDeleteDialog = false;
	}
</script>

<div class="space-y-6">
	<div class="space-y-4">
		<div class="grid">
			<h4 class="mb-2 text-sm font-medium">Export Conversations</h4>

			<p class="mb-4 text-sm text-muted-foreground">
				Download all your conversations as a JSON file. This includes all messages, attachments, and
Repository: ggerganov/llama.cpp
File: ggml/src/ggml-cpu/ggml-cpu.c (C)

Code before:
    return 1;
#else
    return 0;
#endif
}

int ggml_cpu_get_sve_cnt(void) {
#if defined(__ARM_ARCH) && defined(__ARM_FEATURE_SVE)
    return svcntb();
#else
    return 0;
#endif
}

int ggml_cpu_has_sme(void) {
#if defined(__ARM_ARCH) && defined(__ARM_FEATURE_SME)
    return 1;
#else
    return 0;
#endif
}

void ggml_cpu_init(void) {
    // needed to initialize ggml_time
    {
        struct ggml_init_params params = { 0, NULL, false };
        struct ggml_context * ctx = ggml_init(params);
        ggml_free(ctx);
    }

    ggml_critical_section_start();

    static bool is_first_call = true;

    if (is_first_call) {
        // initialize GELU, Quick GELU, SILU and EXP F32 tables
        {
            const uint64_t t_start = ggml_time_us(); UNUSED(t_start);

            for (int i = 0; i < (1 << 16); ++i) {
                union {
                    uint16_t u16;
                    ggml_fp16_t fp16;
                } u = {i};
                float f = GGML_COMPUTE_FP16_TO_FP32(u.fp16);
                ggml_table_f32_f16[i] = f;
                ggml_table_gelu_f16[i] = GGML_CPU_FP32_TO_FP16(ggml_gelu_f32(f));
                ggml_table_gelu_quick_f16[i] = GGML_CPU_FP32_TO_FP16(ggml_gelu_quick_f32(f));
            }

            const uint64_t t_end = ggml_time_us(); UNUSED(t_end);

Review comment (suggestion):
This is a system call that we want to avoid in hot loops. See the discussion starting here: https://github.com/ggml-org/llama.cpp/pull/8709#issuecomment-2252641544

@angt Does this answer the question? I am not very familiar with the Arm feature set logic. Probably it can be improved/extended. I believe the main difficulty is that Arm hardware is not very ubiquitous among regular users (apart from Macs).

Code after:
    return 0;
#endif
}

int ggml_cpu_has_sve(void) {
#if defined(__ARM_ARCH) && defined(__ARM_FEATURE_SVE)
    return 1;
#else
    return 0;
#endif
}

int ggml_cpu_has_matmul_int8(void) {
#if defined(__ARM_ARCH) && defined(__ARM_FEATURE_MATMUL_INT8)
    return 1;
#else
    return 0;
#endif
}

int ggml_cpu_get_sve_cnt(void) {
#if defined(__ARM_ARCH) && defined(__ARM_FEATURE_SVE)
    return ggml_arm_arch_features.sve_cnt;
#else
    return 0;
#endif
}

int ggml_cpu_has_sme(void) {
#if defined(__ARM_ARCH) && defined(__ARM_FEATURE_SME)
    return 1;
#else
    return 0;
#endif
}

void ggml_cpu_init(void) {
    // needed to initialize ggml_time
    {
        struct ggml_init_params params = { 0, NULL, false };
        struct ggml_context * ctx = ggml_init(params);
        ggml_free(ctx);
    }

    ggml_critical_section_start();

    static bool is_first_call = true;

    if (is_first_call) {
        // initialize GELU, Quick GELU, SILU and EXP F32 tables
        {
Repository: ggerganov/llama.cpp
File: tools/server/webui/src/lib/components/app/chat/ChatAttachments/ChatAttachmentThumbnailFile.svelte (Svelte)

Code before:
<script lang="ts">
	import { RemoveButton } from '$lib/components/app';
	import { formatFileSize, getFileTypeLabel, isTextFile } from '$lib/utils';
	import { getPreviewText } from '$lib/utils/text';
	import { AttachmentType } from '$lib/enums';

	interface Props {
		class?: string;
		id: string;
		onClick?: (event?: MouseEvent) => void;
		onRemove?: (id: string) => void;
		name: string;
		readonly?: boolean;
		size?: number;
		textContent?: string;
		// Either uploaded file or stored attachment
		uploadedFile?: ChatUploadedFile;
		attachment?: DatabaseMessageExtra;
	}

	let {
		class: className = '',
		id,
		onClick,
		onRemove,
		name,
		readonly = false,
		size,
		textContent,

Review comment (refactor):
Let's move this import to unified `$lib/utils`

```suggestion
	import { formatFileSize, getFileTypeLabel, getPreviewText, isTextFile } from '$lib/utils';
```

Code after:
<script lang="ts">
	import { RemoveButton } from '$lib/components/app';
	import { formatFileSize, getFileTypeLabel, getPreviewText, isTextFile } from '$lib/utils';
	import { AttachmentType } from '$lib/enums';

	interface Props {
		class?: string;
		id: string;
		onClick?: (event?: MouseEvent) => void;
		onRemove?: (id: string) => void;
		name: string;
		readonly?: boolean;
		size?: number;
		textContent?: string;
		// Either uploaded file or stored attachment
		uploadedFile?: ChatUploadedFile;
		attachment?: DatabaseMessageExtra;
	}

	let {
		class: className = '',
		id,
		onClick,
		onRemove,
		name,
		readonly = false,
		size,
		textContent,
		uploadedFile,
Repository: ggerganov/llama.cpp
File: ggml/src/ggml-cuda/mma.cuh (CUDA)

Code before:
        int32x8_t * acc = (int32x8_t *) D.x;

#if defined(RDNA4)
        using int32x2_t = __attribute__((__vector_size__(2 * sizeof(int)))) int;
        int32x2_t * a_vec = (int32x2_t *) A.x;
        int32x2_t * b_vec = (int32x2_t *) B.x;

        acc[0] = __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32_gfx12(
            true,
            a_vec[0],
            true,
            b_vec[0],
            acc[0],
            true
        );

        acc[0] = __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32_gfx12(
            true,
            a_vec[1],
            true,
            b_vec[1],
            acc[0],
            true
        );

#elif defined(RDNA3) 
        using int32x4_t = __attribute__((__vector_size__(4 * sizeof(int)))) int;
        int32x4_t * a_vec = (int32x4_t *) A.x;
        int32x4_t * b_vec = (int32x4_t *) B.x;

        acc[0] = __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32(
            true,
            a_vec[0],
            true,
            b_vec[0],
            acc[0],
            true
        );

        acc[0] = __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32(
            true,
            a_vec[1],
            true,
            b_vec[1],
            acc[0],
            true
        );
#endif // RDNA4

#else
        GGML_UNUSED_VARS(D, A, B);

Review comment (suggestion):
```suggestion
#elif defined(RDNA3)
```

To fix the EditorConfig CI.

Code after:
        int32x8_t * acc = (int32x8_t *) D.x;

#if defined(RDNA4)
        using int32x2_t = __attribute__((__vector_size__(2 * sizeof(int)))) int;
        int32x2_t * a_vec = (int32x2_t *) A.x;
        int32x2_t * b_vec = (int32x2_t *) B.x;

        acc[0] = __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32_gfx12(
            true,
            a_vec[0],
            true,
            b_vec[0],
            acc[0],
            true
        );

        acc[0] = __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32_gfx12(
            true,
            a_vec[1],
            true,
            b_vec[1],
            acc[0],
            true
        );

#elif defined(RDNA3)
        using int32x4_t = __attribute__((__vector_size__(4 * sizeof(int)))) int;
        int32x4_t * a_vec = (int32x4_t *) A.x;
        int32x4_t * b_vec = (int32x4_t *) B.x;

        acc[0] = __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32(
            true,
            a_vec[0],
            true,
            b_vec[0],
            acc[0],
            true
        );

        acc[0] = __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32(
            true,
            a_vec[1],
            true,
            b_vec[1],
            acc[0],
            true
        );
#endif // RDNA4

#else
        GGML_UNUSED_VARS(D, A, B);
Repository: ggerganov/llama.cpp
File: tools/mtmd/mtmd.h (C/C++)

Code before:
    // return list of pointers to mtmd_bitmap
    // example:
    //   auto bitmaps_c_ptr = bitmaps.c_ptr();
    //   int32_t res = mtmd_tokenize(... bitmaps_c_ptr.data(), bitmaps_c_ptr.size());
    std::vector<const mtmd_bitmap *> c_ptr() {
        std::vector<const mtmd_bitmap *> res(entries.size());
        for (size_t i = 0; i < entries.size(); i++) {
            res[i] = entries[i].ptr.get();
        }
        return res;
    }
};

struct input_chunks {
    input_chunks_ptr ptr;
    input_chunks() = default;
    input_chunks(mtmd_input_chunks * chunks) : ptr(chunks) {}
    ~input_chunks() = default;
    size_t size() { return mtmd_input_chunks_size(ptr.get()); }
    const mtmd_input_chunk * operator[](size_t idx) {
        return mtmd_input_chunks_get(ptr.get(), idx);
    }
};

// insert mtmd_default_marker() into given string, position depends on the projector
std::string mtmd_add_default_marker(mtmd_context *ctx, const std::string &str);

} // namespace mtmd

#endif

#endif


Review comment (suggestion):
I'd prefer to remove API because:
- It's not compatible with pure-C ABI
- The ordering is actually controlled by users via API. This function only change `llama-mtmd-cli`, but make no changes to `llama-server`

Code after:
    // return list of pointers to mtmd_bitmap
    // example:
    //   auto bitmaps_c_ptr = bitmaps.c_ptr();
    //   int32_t res = mtmd_tokenize(... bitmaps_c_ptr.data(), bitmaps_c_ptr.size());
    std::vector<const mtmd_bitmap *> c_ptr() {
        std::vector<const mtmd_bitmap *> res(entries.size());
        for (size_t i = 0; i < entries.size(); i++) {
            res[i] = entries[i].ptr.get();
        }
        return res;
    }
};

struct input_chunks {
    input_chunks_ptr ptr;
    input_chunks() = default;
    input_chunks(mtmd_input_chunks * chunks) : ptr(chunks) {}
    ~input_chunks() = default;
    size_t size() { return mtmd_input_chunks_size(ptr.get()); }
    const mtmd_input_chunk * operator[](size_t idx) {
        return mtmd_input_chunks_get(ptr.get(), idx);
    }
};

} // namespace mtmd

#endif

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

Code before:

    const int64_t i3 = blockIdx.z;
    const int64_t i2 = blockIdx.y;
    const int64_t i1 = blockIdx.x;

    if (i3 >= ne03 || i2 >= ne02 || i1 >= ne01) {
        return;
    }

    const T * src_row = (const T *) ((const char *) src + i1*nb01 + i2*nb02 + i3*nb03);
    T       * dst_row = (T       *) ((      char *) dst + i1*nb1  + i2*nb2  + i3*nb3);

    const int tid = threadIdx.x;
    const int lane_id = tid % WARP_SIZE;

    if (tid >= ne00) {
        return;
    }

    // Phase 1: Each thread processes elements at stride blockDim.x
    // Compute warp-level prefix sums
    for (int64_t i0 = tid; i0 < ne00; i0 += blockDim.x) {
        // Load value and compute prefix sum within warp
        float val = static_cast<float>(src_row[i0]);
        val = warp_prefix_inclusive_sum(val);
        dst_row[i0] = static_cast<T>(val);

        // Last thread of warp stores its sum to shared memory at position based on data index
        if (lane_id == WARP_SIZE - 1 || i0 == ne00 - 1) {
            const int shmem_idx = i0 / WARP_SIZE;
            shmem[shmem_idx] = val;
        }
    }

    // Sync once after all warp prefix sums are computed
    __syncthreads();

    // Phase 2: Add the sum of all preceding warp groups to each element
    for (int64_t i0 = tid; i0 < ne00; i0 += blockDim.x) {
        const int shmem_idx = i0 / WARP_SIZE;
        float sum = 0.0f;
        for (int j = 0; j < shmem_idx; ++j) {
            sum += shmem[j];
        }
        dst_row[i0] = static_cast<T>(static_cast<float>(dst_row[i0]) + sum);
    }
}

template<typename T>
static void cumsum_cuda(
    const T * src, T * dst,

Review comment (suggestion):
It would be much preferable to store the temporary results in registers or shared memory rather than global memory.

Code after:

    __shared__ typename BlockScan::TempStorage temp_storage;
    __shared__ T block_carry;      // carry from previous tile

    const int tid = threadIdx.x;

    const int64_t i1 = blockIdx.x;
    const int64_t i2 = blockIdx.y;
    const int64_t i3 = blockIdx.z;

    if (i1 >= ne01 || i2 >= ne02 || i3 >= ne03) {
        return;
    }

    const T * src_row = src + i1 * s01 + i2 * s02 + i3 * s03;
    T *       dst_row = dst + i1 * s1  + i2 * s2  + i3 * s3;

    if (tid == 0) {
        block_carry = 0;
    }
    __syncthreads();

    for (int64_t start = 0; start < ne00; start += BLOCK_SIZE) {
        int64_t idx = start + tid;
        T x = (idx < ne00) ? src_row[idx] : T(0);

        T inclusive;
        T block_total;
        BlockScan(temp_storage).InclusiveSum(x, inclusive, block_total);

        __syncthreads();

        T final_val = inclusive + block_carry;

        // store result
        if (idx < ne00) {
            dst_row[idx] = final_val;
        }

        __syncthreads();

        if (tid == 0) {
            block_carry += block_total;
        }

        __syncthreads();
    }
#else
    NO_DEVICE_CODE;
#endif // GGML_CUDA_USE_CUB
}
Repository: ggerganov/llama.cpp
File: ggml/src/ggml-cuda/tri.cu (CUDA)

Code before:
#include "tri.cuh"
#include "ggml.h"

// Triangle type comparison - determines which elements to keep
__device__ static inline bool tri_compare(const int i, const int r, const ggml_tri_type type) {
    switch (type) {
        case GGML_TRI_TYPE_LOWER:      return i < r;
        case GGML_TRI_TYPE_LOWER_DIAG: return i <= r;
        case GGML_TRI_TYPE_UPPER:      return i > r;
        case GGML_TRI_TYPE_UPPER_DIAG: return i >= r;
        default: return false;
    }
}

template<typename T>
static __global__ void tri_kernel(
    const T * src, T * dst,
    const int64_t ne00, const int64_t ne01, const int64_t ne02, const int64_t ne03,
    const int64_t nb00, const int64_t nb01, const int64_t nb02, const int64_t nb03,
    const int64_t nb0,  const int64_t nb1,  const int64_t nb2,  const int64_t nb3,
    const ggml_tri_type ttype) {

    const int64_t i3 = blockIdx.z;
    const int64_t i2 = blockIdx.y;
    const int64_t i1 = blockIdx.x;

    if (i3 >= ne03 || i2 >= ne02 || i1 >= ne01) {
        return;
    }

    const T * src_row = (const T *) ((const char *) src + i1*nb01 + i2*nb02 + i3*nb03);
    T       * dst_row = (T       *) ((      char *) dst + i1*nb1  + i2*nb2  + i3*nb3);

    // Each thread processes elements at stride blockDim.x
    for (int64_t i0 = threadIdx.x; i0 < ne00; i0 += blockDim.x) {
        dst_row[i0] = tri_compare(i0, i1, ttype)
            ? src_row[i0] : static_cast<T>(0.f);
    }

Review comment (performance):
This is going to be very slow in GPU code. Preferably make this a constexpr function and provide the `ggml_tri_type` at compile time as a template parameter.

Code after:
#include "common.cuh"
#include "convert.cuh"
#include "tri.cuh"
#include "ggml.h"

template<typename T, bool prefix_keep, int add_to_split>
static __global__ void tri_kernel(
        const T * src, T * dst,
        const int64_t ne00, const int64_t ne01, const int64_t ne02, const int64_t ne03,
        const int64_t nb00, const int64_t nb01, const int64_t nb02, const int64_t nb03,
        const int64_t nb0,  const int64_t nb1,  const int64_t nb2,  const int64_t nb3) {
    const int64_t i3 = blockIdx.z;
    const int64_t i2 = blockIdx.y;
    const int64_t i1 = blockIdx.x;
    const int64_t split_point = i1 + add_to_split;

    GGML_UNUSED_VARS(nb00, nb0);

    if (i3 >= ne03 || i2 >= ne02 || i1 >= ne01) {
        return;
    }

    const T * src_row = src + i1*nb01 + i2*nb02 + i3*nb03;
    T       * dst_row = dst + i1*nb1  + i2*nb2  + i3*nb3;

    if constexpr (prefix_keep) {
        for (int64_t i0 = threadIdx.x; i0 < split_point; i0 += blockDim.x) {
            dst_row[i0] = src_row[i0];
        }
        for (int64_t i0 = threadIdx.x + split_point; i0 < ne00; i0 += blockDim.x) {
            dst_row[i0] = ggml_cuda_cast<T, float>(0.0f);
        }
    } else {
        for (int64_t i0 = threadIdx.x; i0 < split_point; i0 += blockDim.x) {
            dst_row[i0] = ggml_cuda_cast<T, float>(0.0f);
        }
        for (int64_t i0 = threadIdx.x + split_point; i0 < ne00; i0 += blockDim.x) {
            dst_row[i0] = src_row[i0];
Repository: ggerganov/llama.cpp
File: ggml/src/ggml-cuda/cumsum.cu (CUDA)

Code before:
#include <algorithm>
#include "cumsum.cuh"
#include "convert.cuh"
#include "ggml.h"

#ifdef GGML_CUDA_USE_CUB
#   include <cub/device/device_scan.cuh>
#endif


template<typename T, int BLOCK_SIZE>
static __global__ void cumsum_cub_kernel(
    const T* __restrict__ src,
    T* __restrict__ dst,
    const int64_t ne00, const int64_t ne01, const int64_t ne02, const int64_t ne03,
    const int64_t nb01, const int64_t nb02, const int64_t nb03,
    const int64_t nb1,  const int64_t nb2,  const int64_t nb3) {
#ifdef GGML_CUDA_USE_CUB
    using BlockScan = cub::BlockScan<T, BLOCK_SIZE>;

    __shared__ typename BlockScan::TempStorage temp_storage;
    __shared__ T block_carry;      // carry from previous tile

    const int tid = threadIdx.x;

    const int64_t i1 = blockIdx.x;
    const int64_t i2 = blockIdx.y;
    const int64_t i3 = blockIdx.z;

    if (i1 >= ne01 || i2 >= ne02 || i3 >= ne03) {
        return;
    }

    const T* src_row = src + i1 * nb01 + i2 * nb02 + i3 * nb03;
    T*       dst_row = dst + i1 * nb1  + i2 * nb2  + i3 * nb3;

    if (tid == 0) {
        block_carry = 0;
    }
    __syncthreads();

    for (int64_t start = 0; start < ne00; start += BLOCK_SIZE) {

Review comment (suggestion):
```suggestion
        const T * __restrict__ src,
        T * __restrict__ dst,
        const int64_t ne00, const int64_t ne01, const int64_t ne02, const int64_t ne03,
        const int64_t nb01, const int64_t nb02, const int64_t nb03,
        const int64_t nb1,  const int64_t nb2,  const int64_t nb3) {
```

To make it visually distinct from the function body.

Code after:
#include <algorithm>
#include "cumsum.cuh"
#include "convert.cuh"
#include "ggml-cuda/common.cuh"
#include "ggml.h"

#ifdef GGML_CUDA_USE_CUB
#   include <cub/device/device_scan.cuh>
#endif // GGML_CUDA_USE_CUB

template<typename T, int BLOCK_SIZE>
static __global__ void cumsum_cub_kernel(
        const T * __restrict__ src,
        T * __restrict__ dst,
        const int64_t ne00, const int64_t ne01, const int64_t ne02, const int64_t ne03,
        const int64_t  s01, const int64_t  s02, const int64_t  s03,
        const int64_t   s1,  const int64_t   s2,  const int64_t   s3) {
#ifdef GGML_CUDA_USE_CUB
    using BlockScan = cub::BlockScan<T, BLOCK_SIZE>;

    __shared__ typename BlockScan::TempStorage temp_storage;
    __shared__ T block_carry;      // carry from previous tile

    const int tid = threadIdx.x;

    const int64_t i1 = blockIdx.x;
    const int64_t i2 = blockIdx.y;
    const int64_t i3 = blockIdx.z;

    if (i1 >= ne01 || i2 >= ne02 || i3 >= ne03) {
        return;
    }

    const T * src_row = src + i1 * s01 + i2 * s02 + i3 * s03;
    T *       dst_row = dst + i1 * s1  + i2 * s2  + i3 * s3;

    if (tid == 0) {
        block_carry = 0;
    }
    __syncthreads();

    for (int64_t start = 0; start < ne00; start += BLOCK_SIZE) {