Spaces:
Running
Running
ggml : remove oboslete alibi code (skipme) (#0)
Browse files- ggml-cuda/alibi.cu +0 -63
- ggml-cuda/alibi.cuh +0 -5
ggml-cuda/alibi.cu
DELETED
|
@@ -1,63 +0,0 @@
|
|
| 1 |
-
#include "alibi.cuh"
|
| 2 |
-
|
| 3 |
-
static __global__ void alibi_f32(const float * x, float * dst, const int ncols, const int k_rows,
|
| 4 |
-
const int n_heads_log2_floor, const float m0, const float m1) {
|
| 5 |
-
const int col = blockDim.x*blockIdx.x + threadIdx.x;
|
| 6 |
-
|
| 7 |
-
if (col >= ncols) {
|
| 8 |
-
return;
|
| 9 |
-
}
|
| 10 |
-
|
| 11 |
-
const int row = blockDim.y*blockIdx.y + threadIdx.y;
|
| 12 |
-
const int i = row*ncols + col;
|
| 13 |
-
|
| 14 |
-
const int k = row/k_rows;
|
| 15 |
-
|
| 16 |
-
float m_k;
|
| 17 |
-
if (k < n_heads_log2_floor) {
|
| 18 |
-
m_k = powf(m0, k + 1);
|
| 19 |
-
} else {
|
| 20 |
-
m_k = powf(m1, 2 * (k - n_heads_log2_floor) + 1);
|
| 21 |
-
}
|
| 22 |
-
|
| 23 |
-
dst[i] = col * m_k + x[i];
|
| 24 |
-
}
|
| 25 |
-
|
| 26 |
-
static void alibi_f32_cuda(const float * x, float * dst, const int ncols, const int nrows,
|
| 27 |
-
const int k_rows, const int n_heads_log2_floor, const float m0,
|
| 28 |
-
const float m1, cudaStream_t stream) {
|
| 29 |
-
const dim3 block_dims(CUDA_ALIBI_BLOCK_SIZE, 1, 1);
|
| 30 |
-
const int num_blocks_x = (ncols + CUDA_ALIBI_BLOCK_SIZE - 1) / (CUDA_ALIBI_BLOCK_SIZE);
|
| 31 |
-
const dim3 block_nums(num_blocks_x, nrows, 1);
|
| 32 |
-
alibi_f32<<<block_nums, block_dims, 0, stream>>>(x, dst, ncols, k_rows, n_heads_log2_floor, m0, m1);
|
| 33 |
-
}
|
| 34 |
-
|
| 35 |
-
void ggml_cuda_op_alibi(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
|
| 36 |
-
const ggml_tensor * src0 = dst->src[0];
|
| 37 |
-
const float * src0_d = (const float *)src0->data;
|
| 38 |
-
float * dst_d = (float *)dst->data;
|
| 39 |
-
cudaStream_t stream = ctx.stream();
|
| 40 |
-
|
| 41 |
-
GGML_ASSERT(src0->type == GGML_TYPE_F32);
|
| 42 |
-
GGML_ASSERT( dst->type == GGML_TYPE_F32);
|
| 43 |
-
|
| 44 |
-
const int64_t ne00 = src0->ne[0];
|
| 45 |
-
const int64_t ne01 = src0->ne[1];
|
| 46 |
-
const int64_t ne02 = src0->ne[2];
|
| 47 |
-
const int64_t nrows = ggml_nrows(src0);
|
| 48 |
-
|
| 49 |
-
//const int n_past = ((int32_t *) dst->op_params)[0];
|
| 50 |
-
const int n_head = ((int32_t *) dst->op_params)[1];
|
| 51 |
-
float max_bias;
|
| 52 |
-
memcpy(&max_bias, (int32_t *) dst->op_params + 2, sizeof(float));
|
| 53 |
-
|
| 54 |
-
//GGML_ASSERT(ne01 + n_past == ne00);
|
| 55 |
-
GGML_ASSERT(n_head == ne02);
|
| 56 |
-
|
| 57 |
-
const int n_heads_log2_floor = 1 << (int) floor(log2(n_head));
|
| 58 |
-
|
| 59 |
-
const float m0 = powf(2.0f, -(max_bias) / n_heads_log2_floor);
|
| 60 |
-
const float m1 = powf(2.0f, -(max_bias / 2.0f) / n_heads_log2_floor);
|
| 61 |
-
|
| 62 |
-
alibi_f32_cuda(src0_d, dst_d, ne00, nrows, ne01, n_heads_log2_floor, m0, m1, stream);
|
| 63 |
-
}
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
ggml-cuda/alibi.cuh
DELETED
|
@@ -1,5 +0,0 @@
|
|
| 1 |
-
#include "common.cuh"
|
| 2 |
-
|
| 3 |
-
#define CUDA_ALIBI_BLOCK_SIZE 32
|
| 4 |
-
|
| 5 |
-
void ggml_cuda_op_alibi(ggml_backend_cuda_context & ctx, ggml_tensor * dst);
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|