1#include "diagmask.cuh"
2
3static __global__ void diag_mask_inf_f32(const float * x, float * dst, const int ncols, const int rows_per_channel, const int n_past) {
4 const int col = blockDim.y*blockIdx.y + threadIdx.y;
5 const int row = blockDim.x*blockIdx.x + threadIdx.x;
6
7 if (col >= ncols) {
8 return;
9 }
10
11 const int i = row*ncols + col;
12 //dst[i] = col > (n_past + row % rows_per_channel) ? -INFINITY : x[i];
13 //dst[i] = x[i] - (col > n_past + row % rows_per_channel) * INT_MAX; // equivalent within rounding error but slightly faster on GPU
14 dst[i] = x[i] - (col > n_past + row % rows_per_channel) * FLT_MAX;
15}
16
17static void diag_mask_inf_f32_cuda(const float * x, float * dst, const int ncols_x, const int nrows_x, const int rows_per_channel, const int n_past, cudaStream_t stream) {
18 const dim3 block_dims(1, CUDA_DIAG_MASK_INF_BLOCK_SIZE, 1);
19 const int block_num_x = (ncols_x + CUDA_DIAG_MASK_INF_BLOCK_SIZE - 1) / CUDA_DIAG_MASK_INF_BLOCK_SIZE;
20 const dim3 block_nums(nrows_x, block_num_x, 1);
21 diag_mask_inf_f32<<<gridDim: block_nums, blockDim: block_dims, sharedMem: 0, stream>>>(x, dst, ncols: ncols_x, rows_per_channel, n_past);
22}
23
24void ggml_cuda_op_diag_mask_inf(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
25 const ggml_tensor * src0 = dst->src[0];
26 const float * src0_d = (const float *)src0->data;
27 float * dst_d = (float *)dst->data;
28 cudaStream_t stream = ctx.stream();
29
30 GGML_ASSERT(src0->type == GGML_TYPE_F32);
31 GGML_ASSERT( dst->type == GGML_TYPE_F32);
32
33 const int64_t ne00 = src0->ne[0];
34 const int64_t ne01 = src0->ne[1];
35 const int nrows0 = ggml_nrows(src0);
36
37 const int n_past = ((int32_t *) dst->op_params)[0];
38
39 diag_mask_inf_f32_cuda(x: src0_d, dst: dst_d, ncols_x: ne00, nrows_x: nrows0, rows_per_channel: ne01, n_past, stream);
40}
41