Comments 19
Немного не по теме поста. Я на Хабре совсем недавно. Вижу пост в закладки добавлен 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
Первая идея, которая пришла в голову - использовать сетевую дирректорию для ollama. Она по сути использует файлы из папки и на этом всё
Да, так мы сэкономим десятки гигов, но просчеты будут у того, что запустил. Думаю и запуск будет очень долгий из-за заполнения графической памяти из сети, скорость запуска сильно зависит от скорости интерфейса, запуск на мак студии с 128 гигов оперативы быстрее чем на pci5.0x16 ada6000, и в разы быстрее если pci3.0x16 ada6000. Это все касается только LM Studio. Скорость работы, честно не сравнивал. На мак студии студии запучшена более жирная модель. Так как там 80 гигов выделяется под видео память, а у ada6000 48 гигов.
Как запустить свою LLM для инференса. Руководство по запуску: Ollama, vLLM, Triton, LM Studio, llama.cpp, SGLang