Spaces:
Running
Running
Commit
·
310a36c
1
Parent(s):
4aa54ec
SYCL: Fix GGML_SYCL_DEBUG macro (llama/11995)
Browse files
ggml/src/ggml-sycl/common.hpp
CHANGED
|
@@ -35,7 +35,7 @@
|
|
| 35 |
void* ggml_sycl_host_malloc(size_t size);
|
| 36 |
void ggml_sycl_host_free(void* ptr);
|
| 37 |
|
| 38 |
-
|
| 39 |
#define GGML_SYCL_DEBUG(...) \
|
| 40 |
do { \
|
| 41 |
if (g_ggml_sycl_debug) \
|
|
|
|
| 35 |
void* ggml_sycl_host_malloc(size_t size);
|
| 36 |
void ggml_sycl_host_free(void* ptr);
|
| 37 |
|
| 38 |
+
extern int g_ggml_sycl_debug;
|
| 39 |
#define GGML_SYCL_DEBUG(...) \
|
| 40 |
do { \
|
| 41 |
if (g_ggml_sycl_debug) \
|
ggml/src/ggml-sycl/ggml-sycl.cpp
CHANGED
|
@@ -41,6 +41,7 @@
|
|
| 41 |
#include "ggml-sycl/gemm.hpp"
|
| 42 |
|
| 43 |
static bool g_sycl_loaded = false;
|
|
|
|
| 44 |
|
| 45 |
static ggml_sycl_device_info ggml_sycl_init() {
|
| 46 |
ggml_sycl_device_info info = {};
|
|
@@ -157,8 +158,8 @@ static void ggml_check_sycl() try {
|
|
| 157 |
static bool initialized = false;
|
| 158 |
|
| 159 |
if (!initialized) {
|
| 160 |
-
GGML_SYCL_DEBUG("[SYCL] call ggml_check_sycl\n");
|
| 161 |
g_ggml_sycl_debug = get_sycl_env("GGML_SYCL_DEBUG", 0);
|
|
|
|
| 162 |
GGML_LOG_INFO("GGML_SYCL_DEBUG: %d\n", g_ggml_sycl_debug);
|
| 163 |
#if defined(GGML_SYCL_FORCE_MMQ)
|
| 164 |
GGML_LOG_INFO("GGML_SYCL_FORCE_MMQ: yes\n");
|
|
|
|
| 41 |
#include "ggml-sycl/gemm.hpp"
|
| 42 |
|
| 43 |
static bool g_sycl_loaded = false;
|
| 44 |
+
int g_ggml_sycl_debug = 0;
|
| 45 |
|
| 46 |
static ggml_sycl_device_info ggml_sycl_init() {
|
| 47 |
ggml_sycl_device_info info = {};
|
|
|
|
| 158 |
static bool initialized = false;
|
| 159 |
|
| 160 |
if (!initialized) {
|
|
|
|
| 161 |
g_ggml_sycl_debug = get_sycl_env("GGML_SYCL_DEBUG", 0);
|
| 162 |
+
GGML_SYCL_DEBUG("[SYCL] call ggml_check_sycl\n");
|
| 163 |
GGML_LOG_INFO("GGML_SYCL_DEBUG: %d\n", g_ggml_sycl_debug);
|
| 164 |
#if defined(GGML_SYCL_FORCE_MMQ)
|
| 165 |
GGML_LOG_INFO("GGML_SYCL_FORCE_MMQ: yes\n");
|
ggml/src/ggml-sycl/softmax.cpp
CHANGED
|
@@ -249,13 +249,16 @@ void ggml_sycl_op_soft_max(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
|
|
| 249 |
|
| 250 |
if (dst->src[1] && dst->src[1]->type == GGML_TYPE_F16) {
|
| 251 |
const sycl::half * src1_dd = static_cast<sycl::half *>(dst->src[1]->data);
|
|
|
|
| 252 |
soft_max_f32_sycl<sycl::half>(src0_dd, src1_dd, dst_dd, ne00, nrows_x, nrows_y, scale, max_bias,
|
| 253 |
main_stream, ctx.device);
|
| 254 |
} else if (dst->src[1] && dst->src[1]->type == GGML_TYPE_F32) {
|
| 255 |
const float * src1_dd = static_cast<const float *>(dst->src[1]->data);
|
|
|
|
| 256 |
soft_max_f32_sycl<float>(src0_dd, src1_dd, dst_dd, ne00, nrows_x, nrows_y, scale, max_bias, main_stream, ctx.device);
|
| 257 |
} else {
|
| 258 |
/* mask unavailable */
|
|
|
|
| 259 |
soft_max_f32_sycl<float>(src0_dd, nullptr, dst_dd, ne00, nrows_x, nrows_y, scale, max_bias, main_stream, ctx.device);
|
| 260 |
}
|
| 261 |
}
|
|
|
|
| 249 |
|
| 250 |
if (dst->src[1] && dst->src[1]->type == GGML_TYPE_F16) {
|
| 251 |
const sycl::half * src1_dd = static_cast<sycl::half *>(dst->src[1]->data);
|
| 252 |
+
GGML_SYCL_DEBUG("%s: F16 mask\n", __func__);
|
| 253 |
soft_max_f32_sycl<sycl::half>(src0_dd, src1_dd, dst_dd, ne00, nrows_x, nrows_y, scale, max_bias,
|
| 254 |
main_stream, ctx.device);
|
| 255 |
} else if (dst->src[1] && dst->src[1]->type == GGML_TYPE_F32) {
|
| 256 |
const float * src1_dd = static_cast<const float *>(dst->src[1]->data);
|
| 257 |
+
GGML_SYCL_DEBUG("%s: F32 mask\n", __func__);
|
| 258 |
soft_max_f32_sycl<float>(src0_dd, src1_dd, dst_dd, ne00, nrows_x, nrows_y, scale, max_bias, main_stream, ctx.device);
|
| 259 |
} else {
|
| 260 |
/* mask unavailable */
|
| 261 |
+
GGML_SYCL_DEBUG("%s: No mask\n", __func__);
|
| 262 |
soft_max_f32_sycl<float>(src0_dd, nullptr, dst_dd, ne00, nrows_x, nrows_y, scale, max_bias, main_stream, ctx.device);
|
| 263 |
}
|
| 264 |
}
|