1#include "conv-transpose-1d.cuh"
2
3static __global__ void conv_transpose_1d_kernel(
4 const int s0, const int p0, const int d0, const int output_size,
5 const int src0_ne0, const int src0_ne1, const int src0_ne2, const int src0_ne3,
6 const int src1_ne0, const int src1_ne1, const int src1_ne2, const int src1_ne3,
7 const int dst_ne0, const int dst_ne1, const int dst_ne2, const int dst_ne3,
8 const float * src0, const float * src1, float * dst) {
9 int global_index = threadIdx.x + blockIdx.x * blockDim.x;
10 if (global_index >= output_size) {
11 return;
12 }
13
14 int out_index = global_index / dst_ne0;
15
16 float accumulator = 0;
17
18 for (int c = 0; c < src0_ne2; c++) {
19 int idx = global_index % dst_ne0;
20
21 int kernel_offset = (src0_ne0 * src0_ne1 * c) + (out_index * src0_ne0);
22 int input_offset = src1_ne0 * c;
23
24 for (int i = 0; i < src1_ne0; i++) {
25 if (!(idx >= i*s0 && idx < i*s0 + src0_ne0)) {
26 continue;
27 }
28 int weight_idx = idx - i*s0;
29
30 float kernel_weight = src0[kernel_offset + weight_idx];
31 float input_value = src1[input_offset+i];
32
33 accumulator += kernel_weight * input_value;
34 }
35 }
36 dst[global_index] = accumulator;
37 GGML_UNUSED_VARS(p0, d0, src0_ne3, src1_ne3, dst_ne3, src1_ne1, dst_ne1, src1_ne2, dst_ne2);
38}
39
40static void conv_transpose_1d_f32_f32_cuda(
41 const int s0, const int p0, const int d0, const int output_size,
42 const int src0_ne0, const int src0_ne1, const int src0_ne2, const int src0_ne3,
43 const int src1_ne0, const int src1_ne1, const int src1_ne2, const int src1_ne3,
44 const int dst_ne0, const int dst_ne1, const int dst_ne2, const int dst_ne3,
45 const float * src0, const float * src1, float * dst,
46 cudaStream_t stream) {
47
48 const int num_blocks = (output_size + CUDA_CONV_TRANPOSE_1D_BLOCK_SIZE - 1) / CUDA_CONV_TRANPOSE_1D_BLOCK_SIZE;
49 conv_transpose_1d_kernel<<<gridDim: num_blocks,CUDA_CONV_TRANPOSE_1D_BLOCK_SIZE, sharedMem: 0, stream>>>(
50 s0,p0,d0,output_size,
51 src0_ne0, src0_ne1, src0_ne2, src0_ne3,
52 src1_ne0, src1_ne1, src1_ne2, src1_ne3,
53 dst_ne0, dst_ne1, dst_ne2, dst_ne3,
54 src0,src1, dst);
55}
56
57void ggml_cuda_op_conv_transpose_1d(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
58 const ggml_tensor * src0 = dst->src[0];
59 const float * src0_d = (const float *)src0->data;
60
61 const ggml_tensor * src1 = dst->src[1];
62 const float * src1_d = (const float *)src1->data;
63
64 float * dst_d = (float *)dst->data;
65 cudaStream_t stream = ctx.stream();
66
67 GGML_ASSERT(src0->type == GGML_TYPE_F32);
68 GGML_ASSERT( dst->type == GGML_TYPE_F32);
69
70 GGML_ASSERT(ggml_is_contiguous(src0));
71 GGML_ASSERT(ggml_is_contiguous(src1));
72
73 const int32_t * opts = (const int32_t *)dst->op_params;
74
75 const int s0 = opts[0];
76 const int p0 = 0;//opts[3];
77 const int d0 = 1;//opts[4];
78
79 const int64_t output_size = ggml_nelements(dst);
80
81 conv_transpose_1d_f32_f32_cuda(s0, p0, d0, output_size,
82 src0->ne[0], src0->ne[1], src0->ne[2], src0->ne[3],
83 src1->ne[0], src1->ne[1], src1->ne[2], src1->ne[3],
84 dst->ne[0], dst->ne[1], dst->ne[2], dst->ne[3],
85 src0_d, src1_d, dst_d, stream);
86}
87