summaryrefslogtreecommitdiff
path: root/ggml-cuda/scale.cu
diff options
context:
space:
mode:
Diffstat (limited to 'ggml-cuda/scale.cu')
-rw-r--r--ggml-cuda/scale.cu31
1 files changed, 0 insertions, 31 deletions
diff --git a/ggml-cuda/scale.cu b/ggml-cuda/scale.cu
deleted file mode 100644
index 1405e066..00000000
--- a/ggml-cuda/scale.cu
+++ /dev/null
@@ -1,31 +0,0 @@
-#include "scale.cuh"
-
-static __global__ void scale_f32(const float * x, float * dst, const float scale, const int k) {
- const int i = blockDim.x*blockIdx.x + threadIdx.x;
-
- if (i >= k) {
- return;
- }
-
- dst[i] = scale * x[i];
-}
-
-static void scale_f32_cuda(const float * x, float * dst, const float scale, const int k, cudaStream_t stream) {
- const int num_blocks = (k + CUDA_SCALE_BLOCK_SIZE - 1) / CUDA_SCALE_BLOCK_SIZE;
- scale_f32<<<num_blocks, CUDA_SCALE_BLOCK_SIZE, 0, stream>>>(x, dst, scale, k);
-}
-
-void ggml_cuda_op_scale(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
- const ggml_tensor * src0 = dst->src[0];
- const float * src0_d = (const float *)src0->data;
- float * dst_d = (float *)dst->data;
- cudaStream_t stream = ctx.stream();
-
- GGML_ASSERT(src0->type == GGML_TYPE_F32);
- GGML_ASSERT( dst->type == GGML_TYPE_F32);
-
- float scale;
- memcpy(&scale, dst->op_params, sizeof(float));
-
- scale_f32_cuda(src0_d, dst_d, scale, ggml_nelements(src0), stream);
-}