diff options
Diffstat (limited to 'ggml-cuda/scale.cu')
-rw-r--r-- | ggml-cuda/scale.cu | 31 |
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); -} |