forked from ggerganov/llama.cpp
-
Notifications
You must be signed in to change notification settings - Fork 0
Commit
This commit does not belong to any branch on this repository, and may belong to a fork outside of the repository.
[SYCL] add conv support (ggerganov#8688)
- Loading branch information
Showing
5 changed files
with
134 additions
and
0 deletions.
There are no files selected for viewing
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,99 @@ | ||
// | ||
// MIT license | ||
// Copyright (C) 2024 Intel Corporation | ||
// SPDX-License-Identifier: MIT | ||
// | ||
|
||
// | ||
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. | ||
// See https://llvm.org/LICENSE.txt for license information. | ||
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception | ||
// | ||
|
||
#include "conv.hpp" | ||
|
||
static void conv_transpose_1d_kernel( | ||
const int s0, const int output_size, | ||
const int src0_ne0, const int src0_ne1, const int src0_ne2, | ||
const int src1_ne0, const int dst_ne0, | ||
const float * src0, const float * src1, float * dst, | ||
const sycl::nd_item<3> &item_ct1) { | ||
int global_index = item_ct1.get_local_id(2) + | ||
item_ct1.get_group(2) * item_ct1.get_local_range(2); | ||
if (global_index >= output_size) { | ||
return; | ||
} | ||
|
||
int out_index = global_index / dst_ne0; | ||
|
||
float accumulator = 0; | ||
|
||
for (int c = 0; c < src0_ne2; c++) { | ||
int idx = global_index % dst_ne0; | ||
|
||
int kernel_offset = (src0_ne0 * src0_ne1 * c) + (out_index * src0_ne0); | ||
int input_offset = src1_ne0 * c; | ||
|
||
for (int i = 0; i < src1_ne0; i++) { | ||
if (!(idx >= i*s0 && idx < i*s0 + src0_ne0)) { | ||
continue; | ||
} | ||
int weight_idx = idx - i*s0; | ||
|
||
float kernel_weight = src0[kernel_offset + weight_idx]; | ||
float input_value = src1[input_offset+i]; | ||
|
||
accumulator += kernel_weight * input_value; | ||
} | ||
} | ||
dst[global_index] = accumulator; | ||
} | ||
|
||
static void conv_transpose_1d_f32_f32_sycl( | ||
const int s0, const int output_size, | ||
const int src0_ne0, const int src0_ne1, const int src0_ne2, | ||
const int src1_ne0, const int dst_ne0, | ||
const float *src0, const float *src1, float *dst, | ||
const queue_ptr& stream) { | ||
|
||
const int num_blocks = (output_size + SYCL_CONV_TRANPOSE_1D_BLOCK_SIZE - 1) / SYCL_CONV_TRANPOSE_1D_BLOCK_SIZE; | ||
const sycl::range<3> block_dims(1, 1, SYCL_CONV_TRANPOSE_1D_BLOCK_SIZE); | ||
const sycl::range<3> block_nums(1, 1, num_blocks); | ||
stream->parallel_for( | ||
sycl::nd_range<3>( | ||
block_nums * block_dims, block_dims), | ||
[=](sycl::nd_item<3> item_ct1) { | ||
conv_transpose_1d_kernel( | ||
s0, output_size, | ||
src0_ne0, src0_ne1, src0_ne2, | ||
src1_ne0, dst_ne0, | ||
src0, src1, dst, item_ct1); | ||
}); | ||
} | ||
|
||
void ggml_sycl_op_conv_transpose_1d(ggml_backend_sycl_context & ctx, const ggml_tensor *src0, | ||
const ggml_tensor *src1, ggml_tensor *dst) { | ||
const float * src0_d = (const float *)src0->data; | ||
const float * src1_d = (const float *)src1->data; | ||
|
||
float * dst_d = (float *)dst->data; | ||
dpct::queue_ptr stream = ctx.stream(); | ||
|
||
GGML_ASSERT(src0->type == GGML_TYPE_F32); | ||
GGML_ASSERT( dst->type == GGML_TYPE_F32); | ||
|
||
GGML_ASSERT(ggml_is_contiguous(src0)); | ||
GGML_ASSERT(ggml_is_contiguous(src1)); | ||
|
||
const int32_t * opts = (const int32_t *)dst->op_params; | ||
|
||
const int s0 = opts[0]; | ||
|
||
const int64_t output_size = ggml_nelements(dst); | ||
|
||
conv_transpose_1d_f32_f32_sycl(s0, output_size, | ||
src0->ne[0], src0->ne[1], src0->ne[2], | ||
src1->ne[0], dst->ne[0], | ||
src0_d, src1_d, dst_d, stream); | ||
} | ||
|
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,21 @@ | ||
// | ||
// MIT license | ||
// Copyright (C) 2024 Intel Corporation | ||
// SPDX-License-Identifier: MIT | ||
// | ||
|
||
// | ||
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. | ||
// See https://llvm.org/LICENSE.txt for license information. | ||
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception | ||
// | ||
|
||
#ifndef GGML_SYCL_CONV_HPP | ||
#define GGML_SYCL_CONV_HPP | ||
|
||
#include "common.hpp" | ||
|
||
void ggml_sycl_op_conv_transpose_1d(ggml_backend_sycl_context & ctx, const ggml_tensor *src0, | ||
const ggml_tensor *src1, ggml_tensor *dst); | ||
|
||
#endif // GGML_SYCL_CONV_HPP |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters