Comments 17
Немного не по теме поста. Я на Хабре совсем недавно. Вижу пост в закладки добавлен 17 раз и 0 лайков. Здесь на Хабре своя особая атмосфера? Информацию полезную получили, а поблагодарить слабо? Или я что-то не правильно интерепретирую?
Данный движок является бекендом у многих популярных инструментов для инференса моделей. Например у таких как Ollama и LM Studio. Ссылка на движок.
Как-то заинтересовало меня, как люди в таких движках свёртки быстрые на GPU делают. Открыл я вот этот файлик. И что же там я вижу?
template <typename T, typename Layout>
__global__ void conv2d_dw_kernel(const T * __restrict__ input, const T * __restrict__ kernel, T * __restrict__ output,
const int in_w, const int in_h, const int out_w, const int out_h,
const int kernel_w, const int kernel_h, const int stride_x, const int stride_y,
const int padding_x, const int padding_y, const int dilation_x, const int dilation_y,
const int channels, const int batches) {
const int global_idx = blockIdx.x * blockDim.x + threadIdx.x;
const int total_elements = batches * channels * out_h * out_w;
if (global_idx >= total_elements) {
return;
}
conv_params params = { in_w, in_h, out_w, out_h, kernel_w, kernel_h, stride_x,
stride_y, padding_x, padding_y, dilation_x, dilation_y, channels, batches };
int batch_idx, channel_idx, out_y_idx, out_x_idx;
Layout::unpack_indices(global_idx, params, batch_idx, channel_idx, out_y_idx, out_x_idx);
T accumulator = 0;
kernel_bounds bounds = calculate_kernel_bounds(out_x_idx, out_y_idx, params);
for (int kern_y = bounds.y_min; kern_y < bounds.y_max; ++kern_y) {
int in_y_idx = calculate_input_coord(out_y_idx, kern_y, params.stride_y, params.dilation_y, params.padding_y);
for (int kern_x = bounds.x_min; kern_x < bounds.x_max; ++kern_x) {
int in_x_idx = calculate_input_coord(out_x_idx, kern_x, params.stride_x, params.dilation_x, params.padding_x);
const T input_val = input[Layout::input_index(batch_idx, channel_idx, in_y_idx, in_x_idx, params)];
const T kernel_val = kernel[Layout::kernel_index(channel_idx, kern_y, kern_x, params)];
accumulator += input_val * kernel_val;
}
}
output[Layout::output_index(batch_idx, channel_idx, out_y_idx, out_x_idx, params)] = accumulator;
}Это ж свёртка в лоб как на CPU один в один! Такой алгоритм крайне медленный. Гораздо, гораздо (на пару порядков) быстрее на основе умножения матриц с разделяемой памятью. Если вся llama такая, то как же она тормозит-то?
И операция % в CUDA не знаю почему, но нифига не быстрая. Если её заменить на деление и вычитание умноженного значения, быстродействие возрастает.
Ничто не мешает вам переписать этот участок кода и открыть пулреквест на гитхаб llama.cpp, если все так как вы описали, то вам ещё и спасибо скажут
Увы, мешает то, что нужно очень хорошо вникнуть в сам проект. К тому же есть ведь и свои дела и свои проекты. И на всё нужны силы и время.
Они багрепорты на креши со 100% повторяемостью не смотрят, если не выпрашивать, думаете там спасибо кто-то скажет? :))
Вы ещё в Олламу не смотрели...
На одной и той же задаче vLLM в 5 раз быстрее. llama.cpp - где-то посередине.
А вообще, возьня со всеми этими фрейморками напоминает установку драйверов для принтера на юниксе в 90х годах прошлого века.
Загрузил qwen , он полностью влез в видяху. Но при работе почемуто сильно грузит процессор на 59% ровно. А ведь говорят что моделям только видяха нужна. Чтото не так)
Вот повоевал с vllm на rocm, поругался, поматерился и перелез на llama-cpp с vulkan бэкендом. Отказалось, что он стабильнее всего работает, да ещё и быстро.
Короче, бесит vendor lock на nvidia и все попытки amd догнать это через rocm. И ещё в целом низкая инженерная культура MLщиков в плане следования принципу УМВР. Хорошо хоть llama-cpp хоть что-то с этим делают.
Подскажите , если я скачаю модель на компьютер, могу ли я ее обучать на своих данных? Можно ли отключить эту ерунду про моралистику, вида такой вопрос некорректный,на такой отвечать не буду?
LM Studio сегодня меня довел до ручки, и перешел на Ollama, интерфейс сделали очень красивый, я в восторге... oss-120 отлично себя чувствует на мак студио 128Гб, использую удаленно, это значительно дешевле, чем покупать видео карту с таким количество памяти. возьмите на заметку. Если подскажите будет очень классно, Ollama заставить использовать облачные или уделанные модели. Локальную LM Studio можно подключить удаленной LM Studio и использовать в локальной сети всем желающим, хочу повторить данный опыт с Ollama
Как запустить свою LLM для инференса. Руководство по запуску: Ollama, vLLM, Triton, LM Studio, llama.cpp, SGLang