format
This commit is contained in:
parent
b830685378
commit
267af4e75d
1 changed files with 31 additions and 31 deletions
|
@ -63,7 +63,7 @@ static int g_ggml_sycl_debug = 0;
|
||||||
// define for XMX in Intel GPU
|
// define for XMX in Intel GPU
|
||||||
// TODO: currently, it's not used for XMX really.
|
// TODO: currently, it's not used for XMX really.
|
||||||
#if !defined(GGML_SYCL_FORCE_MMQ)
|
#if !defined(GGML_SYCL_FORCE_MMQ)
|
||||||
#define SYCL_USE_XMX
|
#define SYCL_USE_XMX
|
||||||
#endif
|
#endif
|
||||||
|
|
||||||
// max batch size to use MMQ kernels when tensor cores are available
|
// max batch size to use MMQ kernels when tensor cores are available
|
||||||
|
@ -84,16 +84,16 @@ static int g_ggml_sycl_debug = 0;
|
||||||
typedef sycl::queue *queue_ptr;
|
typedef sycl::queue *queue_ptr;
|
||||||
|
|
||||||
enum ggml_sycl_backend_gpu_mode {
|
enum ggml_sycl_backend_gpu_mode {
|
||||||
SYCL_UNSET_GPU_MODE = -1,
|
SYCL_UNSET_GPU_MODE = -1,
|
||||||
SYCL_SINGLE_GPU_MODE = 0,
|
SYCL_SINGLE_GPU_MODE = 0,
|
||||||
SYCL_MUL_GPU_MODE
|
SYCL_MUL_GPU_MODE
|
||||||
};
|
};
|
||||||
|
|
||||||
static_assert(sizeof(sycl::half) == sizeof(ggml_fp16_t), "wrong fp16 size");
|
static_assert(sizeof(sycl::half) == sizeof(ggml_fp16_t), "wrong fp16 size");
|
||||||
|
|
||||||
static void crash() {
|
static void crash() {
|
||||||
int* ptr = NULL;
|
int* ptr = NULL;
|
||||||
*ptr = 0;
|
*ptr = 0;
|
||||||
}
|
}
|
||||||
|
|
||||||
[[noreturn]] static void ggml_sycl_error(
|
[[noreturn]] static void ggml_sycl_error(
|
||||||
|
@ -102,9 +102,9 @@ static void crash() {
|
||||||
const char* file,
|
const char* file,
|
||||||
const int line,
|
const int line,
|
||||||
const char* msg) {
|
const char* msg) {
|
||||||
fprintf(stderr, "SYCL error: %s: %s\n", stmt, msg);
|
fprintf(stderr, "SYCL error: %s: %s\n", stmt, msg);
|
||||||
fprintf(stderr, " in function %s at %s:%d\n", func, file, line);
|
fprintf(stderr, " in function %s at %s:%d\n", func, file, line);
|
||||||
GGML_ABORT("SYCL error");
|
GGML_ABORT("SYCL error");
|
||||||
}
|
}
|
||||||
|
|
||||||
#define SYCL_CHECK(err) \
|
#define SYCL_CHECK(err) \
|
||||||
|
@ -142,40 +142,40 @@ static int g_all_sycl_device_count = -1;
|
||||||
static bool g_ggml_backend_sycl_buffer_type_initialized = false;
|
static bool g_ggml_backend_sycl_buffer_type_initialized = false;
|
||||||
|
|
||||||
static ggml_sycl_backend_gpu_mode g_ggml_sycl_backend_gpu_mode =
|
static ggml_sycl_backend_gpu_mode g_ggml_sycl_backend_gpu_mode =
|
||||||
SYCL_UNSET_GPU_MODE;
|
SYCL_UNSET_GPU_MODE;
|
||||||
|
|
||||||
static void* g_scratch_buffer = nullptr;
|
static void* g_scratch_buffer = nullptr;
|
||||||
static size_t g_scratch_size = 0; // disabled by default
|
static size_t g_scratch_size = 0; // disabled by default
|
||||||
static size_t g_scratch_offset = 0;
|
static size_t g_scratch_offset = 0;
|
||||||
|
|
||||||
[[noreturn]] static inline void bad_arch(const sycl::stream& stream_ct1) {
|
[[noreturn]] static inline void bad_arch(const sycl::stream& stream_ct1) {
|
||||||
stream_ct1 << "ERROR: ggml-sycl was compiled without support for the "
|
stream_ct1 << "ERROR: ggml-sycl was compiled without support for the "
|
||||||
"current GPU architecture.\n";
|
"current GPU architecture.\n";
|
||||||
// __trap();
|
// __trap();
|
||||||
std::exit(1);
|
std::exit(1);
|
||||||
|
|
||||||
(void)bad_arch; // suppress unused function warning
|
(void)bad_arch; // suppress unused function warning
|
||||||
}
|
}
|
||||||
|
|
||||||
int get_current_device_id();
|
int get_current_device_id();
|
||||||
|
|
||||||
inline dpct::err0 ggml_sycl_set_device(const int device) try {
|
inline dpct::err0 ggml_sycl_set_device(const int device) try {
|
||||||
|
|
||||||
int current_device_id;
|
int current_device_id;
|
||||||
SYCL_CHECK(CHECK_TRY_ERROR(current_device_id = get_current_device_id()));
|
SYCL_CHECK(CHECK_TRY_ERROR(current_device_id = get_current_device_id()));
|
||||||
|
|
||||||
// GGML_SYCL_DEBUG("ggml_sycl_set_device device_id=%d,
|
// GGML_SYCL_DEBUG("ggml_sycl_set_device device_id=%d,
|
||||||
// current_device_id=%d\n", device, current_device);
|
// current_device_id=%d\n", device, current_device);
|
||||||
if (device == current_device_id) {
|
if (device == current_device_id) {
|
||||||
return 0;
|
return 0;
|
||||||
}
|
}
|
||||||
|
|
||||||
return CHECK_TRY_ERROR(dpct::select_device(device));
|
return CHECK_TRY_ERROR(dpct::select_device(device));
|
||||||
} catch (sycl::exception const& exc) {
|
} catch (sycl::exception const& exc) {
|
||||||
std::cerr << exc.what() << "Exception caught at file:" << __FILE__
|
std::cerr << exc.what() << "Exception caught at file:" << __FILE__
|
||||||
<< ", line:" << __LINE__ << std::endl;
|
<< ", line:" << __LINE__ << std::endl;
|
||||||
crash();
|
crash();
|
||||||
std::exit(1);
|
std::exit(1);
|
||||||
}
|
}
|
||||||
|
|
||||||
//////////////////////
|
//////////////////////
|
||||||
|
@ -253,10 +253,10 @@ struct ggml_sycl_pool_alloc {
|
||||||
// backend interface
|
// backend interface
|
||||||
|
|
||||||
struct ggml_tensor_extra_gpu {
|
struct ggml_tensor_extra_gpu {
|
||||||
void* data_device[GGML_SYCL_MAX_DEVICES]; // 1 pointer for each device for split
|
void* data_device[GGML_SYCL_MAX_DEVICES]; // 1 pointer for each device for split
|
||||||
// tensors
|
// tensors
|
||||||
dpct::event_ptr events[GGML_SYCL_MAX_DEVICES]
|
dpct::event_ptr events[GGML_SYCL_MAX_DEVICES]
|
||||||
[GGML_SYCL_MAX_STREAMS]; // events for synchronizing multiple GPUs
|
[GGML_SYCL_MAX_STREAMS]; // events for synchronizing multiple GPUs
|
||||||
};
|
};
|
||||||
|
|
||||||
struct ggml_backend_sycl_context {
|
struct ggml_backend_sycl_context {
|
||||||
|
|
Loading…
Add table
Add a link
Reference in a new issue