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<<<block_nums, block_dims, 0, stream>>>(x, dst, 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(src0_d, dst_d, ne00, nrows0, ne01, n_past, stream);
40}