cmake --build buildWithCublas --config Release
```
+**oneDNN**: The current oneDNN releases *(shipped with the oneAPI base-toolkit)* do not include the NVIDIA backend. Therefore, oneDNN must be compiled from source to enable the NVIDIA target:
+
+```sh
+git clone https://github.com/oneapi-src/oneDNN.git
+cd oneDNN
+cmake -GNinja -Bbuild-nvidia -DDNNL_CPU_RUNTIME=DPCPP -DDNNL_GPU_RUNTIME=DPCPP -DDNNL_GPU_VENDOR=NVIDIA -DONEDNN_BUILD_GRAPH=OFF -DCMAKE_C_COMPILER=icx -DCMAKE_CXX_COMPILER=icpx
+cmake --build build-nvidia --config Release
+```
+
- **Adding support to AMD GPUs**
**oneAPI Plugin**: In order to enable SYCL support on AMD GPUs, please install the [Codeplay oneAPI Plugin for AMD GPUs](https://developer.codeplay.com/products/oneapi/amd/download). As with Nvidia GPUs, the user should also make sure the plugin version matches the installed base toolkit.
GGML_SYCL_DEVICE_ARCH=sm_80 # Example architecture
# Option 1: Use FP32 (recommended for better performance in most cases)
-cmake -B build -DGGML_SYCL=ON -DGGML_SYCL_TARGET=NVIDIA -DGGML_SYCL_DEVICE_ARCH=${GGML_SYCL_DEVICE_ARCH} -DCMAKE_C_COMPILER=icx -DCMAKE_CXX_COMPILER=icpx
+cmake -B build -DGGML_SYCL=ON -DGGML_SYCL_TARGET=NVIDIA -DGGML_SYCL_DEVICE_ARCH=${GGML_SYCL_DEVICE_ARCH} -DCMAKE_C_COMPILER=icx -DCMAKE_CXX_COMPILER=icpx -DDNNL_DIR=/path/to/oneDNN/build-nvidia/install/lib/cmake/dnnl
# Option 2: Use FP16
-cmake -B build -DGGML_SYCL=ON -DGGML_SYCL_TARGET=NVIDIA -DGGML_SYCL_DEVICE_ARCH=${GGML_SYCL_DEVICE_ARCH} -DCMAKE_C_COMPILER=icx -DCMAKE_CXX_COMPILER=icpx -DGGML_SYCL_F16=ON
+cmake -B build -DGGML_SYCL=ON -DGGML_SYCL_TARGET=NVIDIA -DGGML_SYCL_DEVICE_ARCH=${GGML_SYCL_DEVICE_ARCH} -DCMAKE_C_COMPILER=icx -DCMAKE_CXX_COMPILER=icpx -DGGML_SYCL_F16=ON -DDNNL_DIR=/path/to/oneDNN/build-nvidia/install/lib/cmake/dnnl
# build all binary
cmake --build build --config Release -j -v
../../include/ggml-sycl.h
)
+find_package(DNNL)
+set(GGML_SYCL_DNNL 0)
+if(DNNL_FOUND)
+ if (DEFINED ENV{ONEAPI_ROOT} AND NOT DEFINED DNNL_GPU_VENDOR)
+ # Assuming oneDNN packaged with oneapi release is used which
+ # supports only intel target
+ set(DNNL_GPU_VENDOR "INTEL")
+ if(NOT "${GGML_SYCL_TARGET}" STREQUAL "INTEL")
+ message(WARNING "oneDNN builds bundled with oneapi release only support INTEL target")
+ endif()
+ endif()
+
+ # Verify oneDNN was compiled for the same target as llama
+ if("${GGML_SYCL_TARGET}" STREQUAL "${DNNL_GPU_VENDOR}")
+ target_link_libraries(ggml-sycl PRIVATE DNNL::dnnl)
+ set(GGML_SYCL_DNNL 1)
+ get_target_property(CONFIGS DNNL::dnnl IMPORTED_CONFIGURATIONS)
+ foreach(CONFIG ${CONFIGS})
+ get_target_property(DNNL_LIB DNNL::dnnl IMPORTED_LOCATION_${CONFIG})
+ message(STATUS "Found oneDNN: ${DNNL_LIB}")
+ endforeach()
+ else()
+ message(WARNING
+ "oneDNN must be compiled for the same target as llama.cpp.
+ llama.cpp: ${GGML_SYCL_TARGET}, oneDNN: ${DNNL_GPU_VENDOR}.
+ Disabling oneDNN support.")
+ endif()
+else()
+ message(STATUS "oneDNN not found, disabling oneDNN support")
+endif()
+target_compile_definitions(ggml-sycl PRIVATE GGML_SYCL_DNNL=${GGML_SYCL_DNNL})
+
if (GGML_SYCL_F16)
if (GGML_SYCL_TARGET STREQUAL "AMD")
message(WARNING "AMD target does not entirely support FP16 in the SYCL backend.")
file(GLOB GGML_SOURCES_SYCL "*.cpp")
target_sources(ggml-sycl PRIVATE ${GGML_HEADERS_SYCL} ${GGML_SOURCES_SYCL})
-find_package(DNNL)
-message("-- DNNL found:" ${DNNL_FOUND})
-
-if (GGML_SYCL_TARGET STREQUAL "INTEL")
- add_compile_definitions(GGML_SYCL_DNNL=${DNNL_FOUND})
-else()
- add_compile_definitions(GGML_SYCL_DNNL=0)
-endif()
-
-if (${DNNL_FOUND} AND GGML_SYCL_TARGET STREQUAL "INTEL")
- target_link_libraries(ggml-sycl PRIVATE DNNL::dnnl)
-endif()
if (WIN32)
find_package(IntelSYCL REQUIRED)
int get_current_device_id();
inline dpct::err0 ggml_sycl_set_device(const int device) try {
-
int current_device_id;
SYCL_CHECK(CHECK_TRY_ERROR(current_device_id = get_current_device_id()));
}
}
+ T * realloc(size_t size) {
+ GGML_ASSERT(pool != nullptr);
+ if (ptr)
+ pool->free(ptr, actual_size);
+ ptr = (T *) pool->alloc(size * sizeof(T), &this->actual_size);
+ return ptr;
+ }
+
// size is in number of elements
T * alloc(size_t size) {
GGML_ASSERT(pool != nullptr);
dnnl::stream stream_dnnl() {
return stream_dnnl(device, 0);
}
+ dnnl::memory get_scratchpad_mem(const dnnl::memory::desc & scratchpad_md,
+ const dnnl::engine & eng, const queue_ptr q) {
+ ggml_sycl_pool_alloc<uint8_t> * pool;
+ auto it = scratchpad_map.find(q);
+ if (it == scratchpad_map.end()) {
+ scratchpad_map[q] = std::make_unique<ggml_sycl_pool_alloc<uint8_t>>(this->pool());
+ pool = scratchpad_map[q].get();
+ } else {
+ pool = it->second.get();
+ }
+
+ size_t scratchpad_size = scratchpad_md.get_size();
+ if (scratchpad_size > pool->actual_size) {
+ pool->realloc(scratchpad_size);
+ }
+ void * mem_ptr = pool->get();
+ return dnnl::memory(scratchpad_md, eng, mem_ptr);
+ }
#endif
// pool
std::unique_ptr<ggml_sycl_pool> pools[GGML_SYCL_MAX_DEVICES];
+ std::unordered_map<sycl::queue *, std::unique_ptr<ggml_sycl_pool_alloc<uint8_t>>> scratchpad_map;
std::unique_ptr<ggml_sycl_pool> host_pools[GGML_SYCL_MAX_DEVICES];
#ifndef GGML_SYCL_GEMM_HPP
#define GGML_SYCL_GEMM_HPP
-#include <fstream>
-#include <iostream>
-
#include "ggml-sycl.h"
#if GGML_SYCL_DNNL
else static_assert(0);
}
- static inline void row_gemm(sycl::queue& q, bool a_trans,
- bool b_trans, int m, int n, int k,
- const void* a, dt at, const void* b, dt bt, void* c, dt ct)
- {
- // Get the device associated with the queue
- sycl::device dev = q.get_device();
- // Get the context associated with the queue
- sycl::context ctx = q.get_context();
- const dnnl::engine eng = dnnl::sycl_interop::make_engine(dev, ctx);
- const dnnl::stream stream = dnnl::sycl_interop::make_stream(eng, q);
+ static inline void row_gemm(ggml_backend_sycl_context & ctx, bool a_trans, bool b_trans, int m, int n, int k,
+ const void * a, dt at, const void * b, dt bt, void * c, dt ct, const queue_ptr & q) {
+ auto stream = ctx.stream_dnnl(q);
+ auto eng = ctx.engine_dnnl(q);
dnnl::memory::dims a_dims = { m, k };
dnnl::memory::dims b_dims = { k, n };
dnnl::memory::dims c_dims = { m, n };
const auto a_in_md = dnnl::memory::desc(a_dims, at, a_trans ? tag::ba : tag::ab);
const auto b_in_md = dnnl::memory::desc(b_dims, bt, b_trans ? tag::ba : tag::ab);
- const auto c_md = dnnl::memory::desc(c_dims, ct, tag::ab);
- auto a_mem = dnnl::memory(a_in_md, eng, const_cast<void*>(a));
- auto b_mem = dnnl::memory(b_in_md, eng, const_cast<void*>(b));
- auto matmul_pd = dnnl::matmul::primitive_desc(eng, a_in_md, b_in_md, c_md);
- auto c_mem = dnnl::memory(matmul_pd.dst_desc(), eng, c);
+ const auto c_md = dnnl::memory::desc(c_dims, ct, tag::ab);
- // Create the primitive.
- auto matmul_prim = dnnl::matmul(matmul_pd);
- // Primitive arguments.
- std::unordered_map<int, dnnl::memory> matmul_args;
- matmul_args.insert({ DNNL_ARG_SRC, a_mem });
- matmul_args.insert({ DNNL_ARG_WEIGHTS, b_mem });
- matmul_args.insert({ DNNL_ARG_DST, c_mem });
+ dnnl::primitive_attr primitive_attr;
+ primitive_attr.set_scratchpad_mode(dnnl::scratchpad_mode::user);
- matmul_prim.execute(stream, matmul_args);
- }
-
-
- static inline void row_gemm(const dnnl::stream& stream, bool a_trans,
- bool b_trans, int m, int n, int k,
- const void* a, dt at, const void* b, dt bt, void* c, dt ct)
- {
- auto const eng = stream.get_engine();
- dnnl::memory::dims a_dims = { m, k };
- dnnl::memory::dims b_dims = { k, n };
- dnnl::memory::dims c_dims = { m, n };
- const auto a_in_md = dnnl::memory::desc(a_dims, at, a_trans ? tag::ba : tag::ab);
- const auto b_in_md = dnnl::memory::desc(b_dims, bt, b_trans ? tag::ba : tag::ab);
- const auto c_md = dnnl::memory::desc(c_dims, ct, tag::ab);
auto a_mem = dnnl::memory(a_in_md, eng, const_cast<void*>(a));
auto b_mem = dnnl::memory(b_in_md, eng, const_cast<void*>(b));
- auto matmul_pd = dnnl::matmul::primitive_desc(eng, a_in_md, b_in_md, c_md);
+ auto matmul_pd = dnnl::matmul::primitive_desc(eng, a_in_md, b_in_md, c_md, primitive_attr);
auto c_mem = dnnl::memory(matmul_pd.dst_desc(), eng, c);
- // Create the primitive.
+ auto scratchpad_md = matmul_pd.scratchpad_desc();
+ auto scratchpad_mem = ctx.get_scratchpad_mem(scratchpad_md, eng, q);
auto matmul_prim = dnnl::matmul(matmul_pd);
- // Primitive arguments.
+
std::unordered_map<int, dnnl::memory> matmul_args;
matmul_args.insert({ DNNL_ARG_SRC, a_mem });
matmul_args.insert({ DNNL_ARG_WEIGHTS, b_mem });
matmul_args.insert({ DNNL_ARG_DST, c_mem });
+ matmul_args.insert({ DNNL_ARG_SCRATCHPAD, scratchpad_mem });
matmul_prim.execute(stream, matmul_args);
}
const to_fp32_sycl_t to_fp32_sycl = ggml_get_to_fp32_sycl(GGML_TYPE_F16, dst);
to_fp32_sycl(dst_f16.get(), dst_dd_i, row_diff*src1_ncols, stream);
#else
- auto dnnl_stream = ctx.stream_dnnl(stream);
- DnnlGemmWrapper::row_gemm(dnnl_stream, false, true, src1_ncols, row_diff, ne10, src1_ptr, DnnlGemmWrapper::to_dt<sycl::half>(),
- src0_ptr, DnnlGemmWrapper::to_dt<sycl::half>(), dst_f16.get(), DnnlGemmWrapper::to_dt<sycl::half>());
+ DnnlGemmWrapper::row_gemm(ctx, false, true, src1_ncols, row_diff, ne10, src1_ptr,
+ DnnlGemmWrapper::to_dt<sycl::half>(), src0_ptr, DnnlGemmWrapper::to_dt<sycl::half>(),
+ dst_f16.get(), DnnlGemmWrapper::to_dt<sycl::half>(), stream);
const to_fp32_sycl_t to_fp32_sycl = ggml_get_to_fp32_sycl(GGML_TYPE_F16, dst);
to_fp32_sycl(dst_f16.get(), dst_dd_i, row_diff* src1_ncols, stream);
#endif
dst_dd_i, ldc)));
# endif
#else
- auto dnnl_stream = ctx.stream_dnnl(stream);
- DnnlGemmWrapper::row_gemm(dnnl_stream, false, true, src1_ncols, row_diff, ne10, src1_ddf1_i, DnnlGemmWrapper::to_dt<float>(),
- src0_ddf_i, DnnlGemmWrapper::to_dt<float>(), dst_dd_i, DnnlGemmWrapper::to_dt<float>());
+ DnnlGemmWrapper::row_gemm(ctx, false, true, src1_ncols, row_diff, ne10, src1_ddf1_i,
+ DnnlGemmWrapper::to_dt<float>(), src0_ddf_i, DnnlGemmWrapper::to_dt<float>(),
+ dst_dd_i, DnnlGemmWrapper::to_dt<float>(), stream);
#endif
}
GGML_UNUSED(dst);
overview / pkg / ggml / sources / llama.cpp / commitdiff