summaryrefslogtreecommitdiff
path: root/llama.cpp/ggml/src/ggml-cuda/binbcast.cu
diff options
context:
space:
mode:
Diffstat (limited to 'llama.cpp/ggml/src/ggml-cuda/binbcast.cu')
-rw-r--r--llama.cpp/ggml/src/ggml-cuda/binbcast.cu504
1 files changed, 504 insertions, 0 deletions
diff --git a/llama.cpp/ggml/src/ggml-cuda/binbcast.cu b/llama.cpp/ggml/src/ggml-cuda/binbcast.cu
new file mode 100644
index 0000000..7339fe0
--- /dev/null
+++ b/llama.cpp/ggml/src/ggml-cuda/binbcast.cu
@@ -0,0 +1,504 @@
+#include "binbcast.cuh"
+#include <cstdint>
+#include <utility>
+
+static __device__ __forceinline__ float op_repeat(const float a, const float b) {
+ return b;
+ GGML_UNUSED(a);
+}
+
+static __device__ __forceinline__ float op_add(const float a, const float b) {
+ return a + b;
+}
+
+static __device__ __forceinline__ float op_sub(const float a, const float b) {
+ return a - b;
+}
+
+static __device__ __forceinline__ float op_mul(const float a, const float b) {
+ return a * b;
+}
+
+static __device__ __forceinline__ float op_div(const float a, const float b) {
+ return a / b;
+}
+
+template <float (*bin_op)(const float, const float),
+ typename src0_t,
+ typename src1_t,
+ typename dst_t,
+ typename... src1_ptrs>
+static __global__ void k_bin_bcast(const src0_t * src0,
+ const src1_t * src1,
+ dst_t * dst,
+ const int ne0,
+ const int ne1,
+ const int ne2,
+ const uint3 ne3,
+ const uint3 ne10,
+ const uint3 ne11,
+ const uint3 ne12,
+ const uint3 ne13,
+ /*const int s0,*/
+ const int s1,
+ const int s2,
+ const int s3,
+ const int s00,
+ const int s01,
+ const int s02,
+ const int s03,
+ const int s10,
+ const int s11,
+ const int s12,
+ const int s13,
+ src1_ptrs... src1s) {
+ const uint32_t i0s = blockDim.x * blockIdx.x + threadIdx.x;
+ const uint32_t i1 = (blockDim.y * blockIdx.y + threadIdx.y);
+ const uint32_t i2 = fastdiv((blockDim.z * blockIdx.z + threadIdx.z), ne3);
+ const uint32_t i3 = (blockDim.z * blockIdx.z + threadIdx.z) - (i2 * ne3.z);
+
+ if (i0s >= (uint32_t)ne0 || i1 >= (uint32_t)ne1 || i2 >= (uint32_t)ne2 || i3 >= ne3.z) {
+ return;
+ }
+
+ const uint32_t i11 = fastmodulo(i1, ne11);
+ const uint32_t i12 = fastmodulo(i2, ne12);
+ const uint32_t i13 = fastmodulo(i3, ne13);
+
+ const size_t i_src0 = i3*s03 + i2*s02 + i1*s01;
+ const size_t i_src1 = i13*s13 + i12*s12 + i11*s11;
+ const size_t i_dst = i3*s3 + i2*s2 + i1*s1;
+
+ const src0_t * src0_row = src0 ? (src0 + i_src0) : nullptr;
+ dst_t * dst_row = dst + i_dst;
+
+ for (int i0 = i0s; i0 < ne0; i0 += blockDim.x * gridDim.x) {
+ const uint32_t i10 = fastmodulo(i0, ne10);
+
+ float result = src0_row ? (float) src0_row[i0*s00] : 0.0f;
+ if constexpr (sizeof...(src1_ptrs) > 0) {
+ result = (..., (result = bin_op(result, (float)src1s[i_src1 + i10*s10])));
+ } else {
+ result = bin_op(result, (float)src1[i_src1 + i10*s10]);
+ }
+
+ dst_row[i0] = (dst_t) result;
+ }
+}
+
+template <float (*bin_op)(const float, const float),
+ typename src0_t,
+ typename src1_t,
+ typename dst_t,
+ typename... src1_ptrs>
+static __global__ void k_bin_bcast_unravel(const src0_t * src0,
+ const src1_t * src1,
+ dst_t * dst,
+ const uint3 ne0,
+ const uint3 ne1,
+ const uint3 ne2,
+ const uint32_t ne3,
+ const uint3 prod_012,
+ const uint3 prod_01,
+ const uint3 ne10,
+ const uint3 ne11,
+ const uint3 ne12,
+ const uint3 ne13,
+ /*const int s0,*/
+ const int s1,
+ const int s2,
+ const int s3,
+ const int s00,
+ const int s01,
+ const int s02,
+ const int s03,
+ const int s10,
+ const int s11,
+ const int s12,
+ const int s13,
+ src1_ptrs... src1s) {
+ const int i = blockDim.x*blockIdx.x + threadIdx.x;
+
+ const uint32_t i3 = fastdiv(i, prod_012);
+ const uint32_t i2 = fastdiv(i - i3 * prod_012.z, prod_01);
+ const uint32_t i1 = fastdiv(i - i3 * prod_012.z - i2 * prod_01.z, ne0);
+ const uint32_t i0 = i - i3 * prod_012.z - i2 * prod_01.z - i1 * ne0.z;
+
+ if (i0 >= ne0.z || i1 >= ne1.z || i2 >= ne2.z || i3 >= ne3) {
+ return;
+ }
+
+ const int i11 = fastmodulo(i1, ne11);
+ const int i12 = fastmodulo(i2, ne12);
+ const int i13 = fastmodulo(i3, ne13);
+
+ const size_t i_src0 = i3*s03 + i2*s02 + i1*s01;
+ const size_t i_src1 = i13*s13 + i12*s12 + i11*s11;
+ const size_t i_dst = i3*s3 + i2*s2 + i1*s1;
+
+ const src0_t * src0_row = src0 ? (src0 + i_src0) : nullptr;
+ dst_t * dst_row = dst + i_dst;
+
+ const int i10 = fastmodulo(i0, ne10);
+
+ float result = src0_row ? (float) src0_row[i0*s00] : 0.0f;
+ if constexpr (sizeof...(src1_ptrs) > 0) {
+ result = (..., (result = bin_op(result, (float)src1s[i_src1 + i10*s10])));
+ } else {
+ result = bin_op(result, (float)src1[i_src1 + i10*s10]);
+ }
+
+ dst_row[i0] = (dst_t) result;
+}
+
+template <float (*bin_op)(const float, const float), typename src0_t, typename src1_t, typename dst_t, size_t... I>
+static void launch_bin_bcast_pack(const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst,
+ const src0_t * src0_dd, const src1_t * src1_dd, dst_t * dst_dd,
+ cudaStream_t stream, std::index_sequence<I...>) {
+ GGML_TENSOR_BINARY_OP_LOCALS
+
+ int nr0 = ne10 / ne0;
+ int nr1 = ne11 / ne1;
+ int nr2 = ne12 / ne2;
+ int nr3 = ne13 / ne3;
+
+ int nr[4] = { nr0, nr1, nr2, nr3 };
+
+ int64_t cne[] = { ne0, ne1, ne2, ne3 };
+ int64_t cne0[] = { ne00, ne01, ne02, ne03 };
+ int64_t cne1[] = { ne10, ne11, ne12, ne13 };
+
+ size_t cnb[] = { nb0, nb1, nb2, nb3 };
+ size_t cnb0[] = { nb00, nb01, nb02, nb03 };
+ size_t cnb1[] = { nb10, nb11, nb12, nb13 };
+
+ auto collapse = [](int64_t cne[]) {
+ cne[0] *= cne[1];
+ cne[1] = cne[2];
+ cne[2] = cne[3];
+ cne[3] = 1;
+ };
+
+ auto collapse_nb = [](size_t cnb[], const int64_t cne[]) {
+ cnb[1] *= cne[1];
+ cnb[2] *= cne[2];
+ cnb[3] *= cne[3];
+ };
+
+ if (ggml_is_contiguous(src0) && ggml_is_contiguous(src1) && !ggml_is_permuted(src0) && !ggml_is_permuted(src1)) {
+ for (int i = 0; i < 4; i++) {
+ if (nr[i] != 1) {
+ break;
+ }
+ if (i > 0) {
+ collapse_nb(cnb, cne);
+ collapse_nb(cnb0, cne0);
+ collapse_nb(cnb1, cne1);
+ collapse(cne);
+ collapse(cne0);
+ collapse(cne1);
+ }
+ }
+ }
+
+ {
+ int64_t ne0 = cne[0];
+ int64_t ne1 = cne[1];
+ int64_t ne2 = cne[2];
+ int64_t ne3 = cne[3];
+
+ //int64_t ne00 = cne0[0]; GGML_UNUSED(ne00);
+ //int64_t ne01 = cne0[1]; GGML_UNUSED(ne01);
+ //int64_t ne02 = cne0[2]; GGML_UNUSED(ne02);
+ //int64_t ne03 = cne0[3]; GGML_UNUSED(ne03);
+
+ size_t nb0 = cnb[0];
+ size_t nb1 = cnb[1];
+ size_t nb2 = cnb[2];
+ size_t nb3 = cnb[3];
+
+ size_t nb00 = cnb0[0];
+ size_t nb01 = cnb0[1];
+ size_t nb02 = cnb0[2];
+ size_t nb03 = cnb0[3];
+
+ size_t nb10 = cnb1[0];
+ size_t nb11 = cnb1[1];
+ size_t nb12 = cnb1[2];
+ size_t nb13 = cnb1[3];
+
+ //size_t s0 = nb0 / sizeof(dst_t);
+ size_t s1 = nb1 / sizeof(dst_t);
+ size_t s2 = nb2 / sizeof(dst_t);
+ size_t s3 = nb3 / sizeof(dst_t);
+
+ size_t s10 = nb10 / sizeof(src1_t);
+ size_t s11 = nb11 / sizeof(src1_t);
+ size_t s12 = nb12 / sizeof(src1_t);
+ size_t s13 = nb13 / sizeof(src1_t);
+
+ size_t s00 = nb00 / sizeof(src0_t);
+ size_t s01 = nb01 / sizeof(src0_t);
+ size_t s02 = nb02 / sizeof(src0_t);
+ size_t s03 = nb03 / sizeof(src0_t);
+
+ GGML_ASSERT(nb0 % sizeof(dst_t) == 0);
+ GGML_ASSERT(nb1 % sizeof(dst_t) == 0);
+ GGML_ASSERT(nb2 % sizeof(dst_t) == 0);
+ GGML_ASSERT(nb3 % sizeof(dst_t) == 0);
+
+ GGML_ASSERT(nb00 % sizeof(src0_t) == 0);
+ GGML_ASSERT(nb01 % sizeof(src0_t) == 0);
+ GGML_ASSERT(nb02 % sizeof(src0_t) == 0);
+ GGML_ASSERT(nb03 % sizeof(src0_t) == 0);
+
+ GGML_ASSERT(nb10 % sizeof(src1_t) == 0);
+ GGML_ASSERT(nb11 % sizeof(src1_t) == 0);
+ GGML_ASSERT(nb12 % sizeof(src1_t) == 0);
+ GGML_ASSERT(nb13 % sizeof(src1_t) == 0);
+
+ const int block_size = 128;
+
+ int64_t hne0 = std::max(ne0 / 2LL, 1LL);
+
+ dim3 block_dims;
+ block_dims.x = std::min<unsigned int>(hne0, block_size);
+ block_dims.y = std::min<unsigned int>(ne1, block_size / block_dims.x);
+ block_dims.z = std::min(std::min<unsigned int>(ne2 * ne3, block_size / block_dims.x / block_dims.y), 64U);
+
+ dim3 block_nums((hne0 + block_dims.x - 1) / block_dims.x, (ne1 + block_dims.y - 1) / block_dims.y,
+ (ne2 * ne3 + block_dims.z - 1) / block_dims.z);
+
+ const uint3 ne10 = init_fastdiv_values((uint32_t) cne1[0]);
+ const uint3 ne11 = init_fastdiv_values((uint32_t) cne1[1]);
+ const uint3 ne12 = init_fastdiv_values((uint32_t) cne1[2]);
+ const uint3 ne13 = init_fastdiv_values((uint32_t) cne1[3]);
+
+ if (block_nums.z > 65535 || block_nums.y > 65535) {
+ int block_num = (ne0 * ne1 * ne2 * ne3 + block_size - 1) / block_size;
+ const uint3 prod_012 = init_fastdiv_values((uint32_t) (ne0 * ne1 * ne2));
+ const uint3 prod_01 = init_fastdiv_values((uint32_t) (ne0 * ne1));
+ const uint3 ne0_fastdiv = init_fastdiv_values((uint32_t) ne0);
+ const uint3 ne1_fastdiv = init_fastdiv_values((uint32_t) ne1);
+ const uint3 ne2_fastdiv = init_fastdiv_values((uint32_t) ne2);
+
+ if constexpr (sizeof...(I) > 0) {
+ k_bin_bcast_unravel<bin_op, src0_t, src1_t, dst_t><<<block_num, block_size, 0, stream>>>(
+ src0_dd, src1_dd, dst_dd, ne0_fastdiv, ne1_fastdiv, ne2_fastdiv, ne3, prod_012, prod_01, ne10, ne11,
+ ne12, ne13,
+ /*s0,*/ s1, s2, s3,
+ s00, s01, s02, s03,
+ s10, s11, s12, s13, (const src1_t *) dst->src[I + 1]->data...);
+ } else {
+ k_bin_bcast_unravel<bin_op, src0_t, src1_t, dst_t>
+ <<<block_num, block_size, 0, stream>>>(src0_dd, src1_dd, dst_dd, ne0_fastdiv, ne1_fastdiv,
+ ne2_fastdiv, ne3, prod_012, prod_01, ne10, ne11, ne12, ne13,
+ /*s0,*/ s1, s2, s3,
+ s00, s01, s02, s03,
+ s10, s11, s12, s13);
+ }
+ } else {
+ const uint3 ne3_fastdiv = init_fastdiv_values((uint32_t) ne3);
+ if constexpr (sizeof...(I) > 0) {
+ k_bin_bcast<bin_op, src0_t, src1_t, dst_t><<<block_nums, block_dims, 0, stream>>>(
+ src0_dd, src1_dd, dst_dd, ne0, ne1, ne2, ne3_fastdiv, ne10, ne11, ne12, ne13,
+ /*s0,*/ s1, s2, s3,
+ s00 ,s01, s02, s03,
+ s10, s11, s12, s13, (const src1_t *) dst->src[I + 1]->data...);
+ } else {
+ k_bin_bcast<bin_op, src0_t, src1_t, dst_t><<<block_nums, block_dims, 0, stream>>>(
+ src0_dd, src1_dd, dst_dd, ne0, ne1, ne2, ne3_fastdiv, ne10, ne11, ne12, ne13,
+ /*s0,*/ s1, s2, s3,
+ s00, s01, s02, s03,
+ s10, s11, s12, s13);
+ }
+ }
+ }
+}
+
+template <typename T>
+static __global__ void k_repeat_back(
+ const T * __restrict__ src, T * __restrict__ dst, const int64_t ne00, const int64_t ne01, const int64_t ne02, const int64_t ne03,
+ const size_t s00, const size_t s01, const size_t s02, const size_t s03,
+ const int64_t ne0, const int64_t ne1, const int64_t ne2, const int64_t ne3) {
+
+ const int64_t tid0 = int64_t(blockIdx.x)*blockDim.x + threadIdx.x;
+ const int64_t tid1 = int64_t(blockIdx.y)*blockDim.y + threadIdx.y;
+ const int64_t tid23 = int64_t(blockIdx.z)*blockDim.z + threadIdx.z;
+ const int64_t tid2 = tid23 % ne2;
+ const int64_t tid3 = tid23 / ne2;
+
+ if (tid0 >= ne0) {
+ return;
+ }
+
+ T sum = 0;
+ for (int64_t i3 = tid3; i3 < ne03; i3 += ne3) {
+ for (int64_t i2 = tid2; i2 < ne02; i2 += ne2) {
+ for (int64_t i1 = tid1; i1 < ne01; i1 += ne1) {
+ for (int64_t i0 = tid0; i0 < ne00; i0 += ne0) {
+ sum += src[i3*s03 + i2*s02 + i1*s01 + i0*s00];
+ }
+ }
+ }
+ }
+ dst[tid3*ne2*ne1*ne0 + tid2*ne1*ne0 + tid1*ne0 + tid0] = sum;
+}
+
+template <float (*bin_op)(const float, const float), int n_fuse = 1>
+struct bin_bcast_cuda {
+ template<typename src0_t, typename src1_t, typename dst_t>
+ void operator()(const struct ggml_tensor * src0, const struct ggml_tensor * src1, struct ggml_tensor * dst,
+ const src0_t * src0_dd, const src1_t * src1_dd, dst_t * dst_dd,
+ cudaStream_t stream) {
+ launch_bin_bcast_pack<bin_op, src0_t, src1_t, dst_t>(
+ src0, src1, dst, src0_dd, src1_dd, dst_dd, stream, std::make_index_sequence<n_fuse>{});
+ }
+};
+
+template <typename T>
+static void repeat_back_cuda(
+ const T * src, T * dst, const int64_t ne00, const int64_t ne01, const int64_t ne02, const int64_t ne03,
+ const size_t s00, const size_t s01, const size_t s02, const size_t s03,
+ const int64_t ne0, const int64_t ne1, const int64_t ne2, const int64_t ne3, cudaStream_t stream) {
+
+ const dim3 block_dims(WARP_SIZE, 1, 1);
+ const dim3 block_nums((ne0 + WARP_SIZE - 1) / WARP_SIZE, ne1, ne2*ne3);
+ k_repeat_back<T><<<block_nums, block_dims, 0, stream>>>
+ (src, dst, ne00, ne01, ne02, ne03, s00, s01, s02, s03, ne0, ne1, ne2, ne3);
+}
+
+template<class op>
+static void ggml_cuda_op_bin_bcast(
+ const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst,
+ const void * src0_dd, const void * src1_dd, void * dst_dd, cudaStream_t stream) {
+
+ GGML_ASSERT(src1->type == GGML_TYPE_F32 || src1->type == GGML_TYPE_F16);
+
+ if (src0->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F32) {
+ op()(src0, src1, dst, (const float *)src0_dd, (const float *)src1_dd, (float *)dst_dd, stream);
+ } else if (src0->type == GGML_TYPE_F16 && src1->type == GGML_TYPE_F16 && dst->type == GGML_TYPE_F16) {
+ op()(src0, src1, dst, (const half *) src0_dd, (const half *)src1_dd, (half *) dst_dd, stream);
+ } else if (src0->type == GGML_TYPE_F16 && src1->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F16) {
+ op()(src0, src1, dst, (const half *) src0_dd, (const float *)src1_dd, (half *) dst_dd, stream);
+ } else if (src0->type == GGML_TYPE_F16 && dst->type == GGML_TYPE_F32) {
+ op()(src0, src1, dst, (const half *) src0_dd, (const float *)src1_dd, (float *)dst_dd, stream);
+ } else {
+ fprintf(stderr, "%s: unsupported types: dst: %s, src0: %s, src1: %s\n", __func__,
+ ggml_type_name(dst->type), ggml_type_name(src0->type), ggml_type_name(src1->type));
+ GGML_ABORT("fatal error");
+ }
+}
+
+void ggml_cuda_op_repeat(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
+ ggml_cuda_op_bin_bcast<bin_bcast_cuda<op_repeat, 0>>(dst, dst->src[0], dst, nullptr, dst->src[0]->data, dst->data, ctx.stream());
+}
+
+void ggml_cuda_op_add(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
+ ggml_cuda_op_bin_bcast<bin_bcast_cuda<op_add>>(dst->src[0], dst->src[1], dst, dst->src[0]->data, dst->src[1]->data, dst->data, ctx.stream());
+}
+
+void ggml_cuda_op_sub(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
+ ggml_cuda_op_bin_bcast<bin_bcast_cuda<op_sub>>(dst->src[0], dst->src[1], dst, dst->src[0]->data, dst->src[1]->data, dst->data, ctx.stream());
+}
+
+void ggml_cuda_op_mul(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
+ ggml_cuda_op_bin_bcast<bin_bcast_cuda<op_mul>>(dst->src[0], dst->src[1], dst, dst->src[0]->data, dst->src[1]->data, dst->data, ctx.stream());
+}
+
+void ggml_cuda_op_div(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
+ ggml_cuda_op_bin_bcast<bin_bcast_cuda<op_div>>(dst->src[0], dst->src[1], dst, dst->src[0]->data, dst->src[1]->data, dst->data, ctx.stream());
+}
+
+template <float (*op)(const float, const float), int n_fuse>
+static void ggml_cuda_op_fused_binbcast_impl(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
+ cudaStream_t stream = ctx.stream();
+
+ const ggml_tensor * src0 = dst->src[0];
+ const ggml_tensor * src1 = dst->src[1];
+
+ if (src0->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F32) {
+ launch_bin_bcast_pack<op, float, float, float>(src0, src1, dst,
+ (const float *) src0->data, (const float *) src1->data, (float *) dst->data,
+ stream, std::make_index_sequence<n_fuse>{});
+ } else if (src0->type == GGML_TYPE_F16 && src1->type == GGML_TYPE_F16 && dst->type == GGML_TYPE_F16) {
+ launch_bin_bcast_pack<op, half, half, half>(src0, src1, dst,
+ (const half *) src0->data, (const half *) src1->data, (half *) dst->data,
+ stream, std::make_index_sequence<n_fuse>{});
+ } else if (src0->type == GGML_TYPE_F16 && src1->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F16) {
+ launch_bin_bcast_pack<op, half, float, half>(src0, src1, dst,
+ (const half *) src0->data, (const float *) src1->data, (half *) dst->data,
+ stream, std::make_index_sequence<n_fuse>{});
+ } else if (src0->type == GGML_TYPE_F16 && dst->type == GGML_TYPE_F32) {
+ launch_bin_bcast_pack<op, half, float, float>(src0, src1, dst,
+ (const half *) src0->data, (const float *) src1->data, (float *) dst->data,
+ stream, std::make_index_sequence<n_fuse>{});
+ } else {
+ fprintf(stderr,
+ "%s: unsupported types for fusion: dst: %s, src0: %s, src1: %s\n",
+ __func__, ggml_type_name(dst->type), ggml_type_name(src0->type), ggml_type_name(src1->type));
+ GGML_ABORT("fatal error");
+ }
+}
+
+
+void ggml_cuda_op_fused_add(ggml_backend_cuda_context & ctx, ggml_tensor * dst, int n_fuse) {
+ GGML_ASSERT(2 <= n_fuse && n_fuse <= 8);
+
+ switch (n_fuse) {
+ case 2:
+ ggml_cuda_op_fused_binbcast_impl<op_add, 2>(ctx, dst);
+ break;
+ case 3:
+ ggml_cuda_op_fused_binbcast_impl<op_add, 3>(ctx, dst);
+ break;
+ case 4:
+ ggml_cuda_op_fused_binbcast_impl<op_add, 4>(ctx, dst);
+ break;
+ case 5:
+ ggml_cuda_op_fused_binbcast_impl<op_add, 5>(ctx, dst);
+ break;
+ case 6:
+ ggml_cuda_op_fused_binbcast_impl<op_add, 6>(ctx, dst);
+ break;
+ case 7:
+ ggml_cuda_op_fused_binbcast_impl<op_add, 7>(ctx, dst);
+ break;
+ case 8:
+ ggml_cuda_op_fused_binbcast_impl<op_add, 8>(ctx, dst);
+ break;
+ default:
+ GGML_ASSERT(false && "Unsupported n_fuse value");
+ }
+}
+
+void ggml_cuda_op_repeat_back(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
+ const ggml_tensor * src0 = dst->src[0];
+
+ GGML_ASSERT(src0->type == dst->type);
+ GGML_ASSERT(ggml_is_contiguous(dst));
+ GGML_ASSERT(ggml_can_repeat(dst, src0));
+
+ cudaStream_t stream = ctx.stream();
+
+ GGML_TENSOR_UNARY_OP_LOCALS;
+
+ GGML_ASSERT(ne2*ne3 <= (1 << 15));
+
+ const size_t ts = ggml_type_size(src0->type);
+ const size_t s00 = nb00 / ts;
+ const size_t s01 = nb01 / ts;
+ const size_t s02 = nb02 / ts;
+ const size_t s03 = nb03 / ts;
+
+ switch (dst->type) {
+ case GGML_TYPE_F32: {
+ const float * src0_d = (const float *) src0->data;
+ float * dst_d = (float *) dst->data;
+ repeat_back_cuda(src0_d, dst_d, ne00, ne01, ne02, ne03, s00, s01, s02, s03, ne0, ne1, ne2, ne3, stream);
+ } break;
+ default: {
+ GGML_ASSERT(false);
+ } break;
+ }
+}
generated by cgit v1.2.3 (git 2.46.0) at 2026-04-29 12:16:24 +0000