89 lines
3.3 KiB
Plaintext
89 lines
3.3 KiB
Plaintext
#include "conv-transpose-1d.cuh"
|
|
|
|
static __global__ void conv_transpose_1d_kernel(
|
|
const int s0, const int p0, const int d0, const int output_size,
|
|
const int src0_ne0, const int src0_ne1, const int src0_ne2, const int src0_ne3,
|
|
const int src1_ne0, const int src1_ne1, const int src1_ne2, const int src1_ne3,
|
|
const int dst_ne0, const int dst_ne1, const int dst_ne2, const int dst_ne3,
|
|
const float * src0, const float * src1, float * dst) {
|
|
int global_index = threadIdx.x + blockIdx.x * blockDim.x;
|
|
if (global_index >= output_size) {
|
|
return;
|
|
}
|
|
|
|
int out_t = global_index % dst_ne0;
|
|
int out_ch = (global_index / dst_ne0) % dst_ne1;
|
|
int plane = global_index / (dst_ne0 * dst_ne1);
|
|
|
|
float accumulator = 0;
|
|
|
|
for (int c = 0; c < src0_ne2; c++) {
|
|
int kernel_offset = src0_ne0 * (out_ch + src0_ne1 * c);
|
|
int input_offset = src1_ne0 * (c + src1_ne1 * plane);
|
|
|
|
for (int k = 0; k < src0_ne0; k++) {
|
|
int input_numer = out_t + p0 - k*d0;
|
|
if (input_numer < 0 || input_numer % s0 != 0) {
|
|
continue;
|
|
}
|
|
|
|
int input_t = input_numer / s0;
|
|
if (input_t >= src1_ne0) {
|
|
continue;
|
|
}
|
|
|
|
accumulator += src0[kernel_offset + k] * src1[input_offset + input_t];
|
|
}
|
|
}
|
|
dst[global_index] = accumulator;
|
|
GGML_UNUSED_VARS(src0_ne3, src1_ne2, src1_ne3, dst_ne2, dst_ne3);
|
|
}
|
|
|
|
static void conv_transpose_1d_f32_f32_cuda(
|
|
const int s0, const int p0, const int d0, const int output_size,
|
|
const int src0_ne0, const int src0_ne1, const int src0_ne2, const int src0_ne3,
|
|
const int src1_ne0, const int src1_ne1, const int src1_ne2, const int src1_ne3,
|
|
const int dst_ne0, const int dst_ne1, const int dst_ne2, const int dst_ne3,
|
|
const float * src0, const float * src1, float * dst,
|
|
cudaStream_t stream) {
|
|
|
|
const int num_blocks = (output_size + CUDA_CONV_TRANPOSE_1D_BLOCK_SIZE - 1) / CUDA_CONV_TRANPOSE_1D_BLOCK_SIZE;
|
|
conv_transpose_1d_kernel<<<num_blocks,CUDA_CONV_TRANPOSE_1D_BLOCK_SIZE, 0, stream>>>(
|
|
s0,p0,d0,output_size,
|
|
src0_ne0, src0_ne1, src0_ne2, src0_ne3,
|
|
src1_ne0, src1_ne1, src1_ne2, src1_ne3,
|
|
dst_ne0, dst_ne1, dst_ne2, dst_ne3,
|
|
src0,src1, dst);
|
|
}
|
|
|
|
void ggml_cuda_op_conv_transpose_1d(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
|
|
const ggml_tensor * src0 = dst->src[0];
|
|
const float * src0_d = (const float *)src0->data;
|
|
|
|
const ggml_tensor * src1 = dst->src[1];
|
|
const float * src1_d = (const float *)src1->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);
|
|
|
|
GGML_ASSERT(ggml_is_contiguous(src0));
|
|
GGML_ASSERT(ggml_is_contiguous(src1));
|
|
|
|
const int32_t * opts = (const int32_t *)dst->op_params;
|
|
|
|
const int s0 = opts[0];
|
|
const int p0 = 0;//opts[3];
|
|
const int d0 = 1;//opts[4];
|
|
|
|
const int64_t output_size = ggml_nelements(dst);
|
|
|
|
conv_transpose_1d_f32_f32_cuda(s0, p0, d0, output_size,
|
|
src0->ne[0], src0->ne[1], src0->ne[2], src0->ne[3],
|
|
src1->ne[0], src1->ne[1], src1->ne[2], src1->ne[3],
|
|
dst->ne[0], dst->ne[1], dst->ne[2], dst->ne[3],
|
|
src0_d, src1_d, dst_d, stream);
|
|
}
|