Hi Philip, Am Fr., 9. Okt. 2020 um 18:33 Uhr schrieb Philip Langdale <phil...@overt.org >:
> I've been sitting on this for a couple of years now, and I figured I > should just send it out. This is what I believe is a conceptually > correct port of bwdif to cuda (modulo edge handling which is not done > in the same way because the conditional checks for edges are expensive > in cuda, but that's the same as for yadif_cuda). > > However, I see glitches in some samples where black or white pixels > appear in white or black areas respectively. This seems like some > sort of under/overflow. I've tried to use the largest cuda types > everywhere, and that did appear to improve things but didn't make > it go away. This is what led to me never sending this diff over the > years, but maybe someone else has insights about this. > I am not familiar with cuda. So here is just one difference, which I noticed compared to the c code. Maybe that is the reason for the glitches. > --- > configure | 2 + > libavfilter/Makefile | 2 + > libavfilter/allfilters.c | 1 + > libavfilter/vf_bwdif_cuda.c | 394 +++++++++++++++++++++++++++++++++++ > libavfilter/vf_bwdif_cuda.cu | 290 ++++++++++++++++++++++++++ > 5 files changed, 689 insertions(+) > create mode 100644 libavfilter/vf_bwdif_cuda.c > create mode 100644 libavfilter/vf_bwdif_cuda.cu > > ... > + > +template<typename T> > +__inline__ __device__ T filter(T A, T B, T C, T D, > + T a, T b, T c, T d, T e, T f, T g, > + T h, T i, T j, T k, T l, T m, T n, > + int clip_max) > +{ > + T final; > + > + int fc = C; > + int fd = (c + l) >> 1; > + int fe = B; > In the following you sometimes use B and C directly and sometimes fc and fe. Is there a reason for this? > + > + int temporal_diff0 = abs(c - l); > + int temporal_diff1 = (abs(g - fc) + abs(f - fe)) >> 1; > + int temporal_diff2 = (abs(i - fc) + abs(h - fe)) >> 1; > + int diff = max3(temporal_diff0 >> 1, temporal_diff1, temporal_diff2); > + > + if (!diff) { > + final = fd; > + } else { > + int fb = ((d + m) >> 1) - fc; > + int ff = ((c + l) >> 1) - fe; > If I don“t miss anything this should be: int ff = ((b + k) >> 1) - fe; > + int dc = fd - fc; > + int de = fd - fe; > + int mmax = max3(de, dc, min(fb, ff)); > + int mmin = min3(de, dc, max(fb, ff)); > + diff = max3(diff, mmin, -mmax); > + > + int interpol; > + if (abs(fc - fe) > temporal_diff0) { > + interpol = (((coef_hf[0] * (c + l) > + - coef_hf[1] * (d + m + b + k) > + + coef_hf[2] * (e + n + a + j)) >> 2) > + + coef_lf[0] * (C + B) - coef_lf[1] * (D + A)) >> 13; > + } else { > + interpol = (coef_sp[0] * (C + B) - coef_sp[1] * (D + A)) >> > 13; > + } > + if (interpol > fd + diff) { > + interpol = fd + diff; > + } else if (interpol < fd - diff) { > + interpol = fd - diff; > + } > + final = clip(interpol, 0, clip_max); > + } > + > + return final; > +} > + > +template<typename T> > +__inline__ __device__ void bwdif_single(T *dst, > + cudaTextureObject_t prev, > + cudaTextureObject_t cur, > + cudaTextureObject_t next, > + int dst_width, int dst_height, > int dst_pitch, > + int src_width, int src_height, > + int parity, int tff, bool > skip_spatial_check, > + int clip_max) > +{ > + // Identify location > + int xo = blockIdx.x * blockDim.x + threadIdx.x; > + int yo = blockIdx.y * blockDim.y + threadIdx.y; > + > + if (xo >= dst_width || yo >= dst_height) { > + return; > + } > + > + // Don't modify the primary field > + if (yo % 2 == parity) { > + dst[yo*dst_pitch+xo] = tex2D<T>(cur, xo, yo); > + return; > + } > + > + T A = tex2D<T>(cur, xo, yo + 3); > + T B = tex2D<T>(cur, xo, yo + 1); > + T C = tex2D<T>(cur, xo, yo - 1); > + T D = tex2D<T>(cur, xo, yo - 3); > + > + // Calculate temporal prediction > + int is_second_field = !(parity ^ tff); > + > + cudaTextureObject_t prev2 = prev; > + cudaTextureObject_t prev1 = is_second_field ? cur : prev; > + cudaTextureObject_t next1 = is_second_field ? next : cur; > + cudaTextureObject_t next2 = next; > + > + T a = tex2D<T>(prev2, xo, yo + 4); > + T b = tex2D<T>(prev2, xo, yo + 2); > + T c = tex2D<T>(prev2, xo, yo + 0); > + T d = tex2D<T>(prev2, xo, yo - 2); > + T e = tex2D<T>(prev2, xo, yo - 4); > + T f = tex2D<T>(prev1, xo, yo + 1); > + T g = tex2D<T>(prev1, xo, yo - 1); > + T h = tex2D<T>(next1, xo, yo + 1); > + T i = tex2D<T>(next1, xo, yo - 1); > + T j = tex2D<T>(next2, xo, yo + 4); > + T k = tex2D<T>(next2, xo, yo + 2); > + T l = tex2D<T>(next2, xo, yo + 0); > + T m = tex2D<T>(next2, xo, yo - 2); > + T n = tex2D<T>(next2, xo, yo - 4); > + > + dst[yo*dst_pitch+xo] = filter(A, B, C, D, > + a, b, c, d, e, f, g, > + h, i, j, k, l, m, n, > + clip_max); > +} > + > +template <typename T> > +__inline__ __device__ void bwdif_double(T *dst, > + cudaTextureObject_t prev, > + cudaTextureObject_t cur, > + cudaTextureObject_t next, > + int dst_width, int dst_height, > int dst_pitch, > + int src_width, int src_height, > + int parity, int tff, bool > skip_spatial_check, > + int clip_max) > +{ > + int xo = blockIdx.x * blockDim.x + threadIdx.x; > + int yo = blockIdx.y * blockDim.y + threadIdx.y; > + > + if (xo >= dst_width || yo >= dst_height) { > + return; > + } > + > + if (yo % 2 == parity) { > + // Don't modify the primary field > + dst[yo*dst_pitch+xo] = tex2D<T>(cur, xo, yo); > + return; > + } > + > + T A = tex2D<T>(cur, xo, yo + 3); > + T B = tex2D<T>(cur, xo, yo + 1); > + T C = tex2D<T>(cur, xo, yo - 1); > + T D = tex2D<T>(cur, xo, yo - 3); > + > + // Calculate temporal prediction > + int is_second_field = !(parity ^ tff); > + > + cudaTextureObject_t prev2 = prev; > + cudaTextureObject_t prev1 = is_second_field ? cur : prev; > + cudaTextureObject_t next1 = is_second_field ? next : cur; > + cudaTextureObject_t next2 = next; > + > + T a = tex2D<T>(prev2, xo, yo + 4); > + T b = tex2D<T>(prev2, xo, yo + 2); > + T c = tex2D<T>(prev2, xo, yo + 0); > + T d = tex2D<T>(prev2, xo, yo - 2); > + T e = tex2D<T>(prev2, xo, yo - 4); > + T f = tex2D<T>(prev1, xo, yo + 1); > + T g = tex2D<T>(prev1, xo, yo - 1); > + T h = tex2D<T>(next1, xo, yo + 1); > + T i = tex2D<T>(next1, xo, yo - 1); > + T j = tex2D<T>(next2, xo, yo + 4); > + T k = tex2D<T>(next2, xo, yo + 2); > + T l = tex2D<T>(next2, xo, yo + 0); > + T m = tex2D<T>(next2, xo, yo - 2); > + T n = tex2D<T>(next2, xo, yo - 4); > + > + T final; > + final.x = filter(A.x, B.x, C.x, D.x, > + a.x, b.x, c.x, d.x, e.x, f.x, g.x, > + h.x, i.x, j.x, k.x, l.x, m.x, n.x, > + clip_max); > + final.y = filter(A.y, B.y, C.y, D.y, > + a.y, b.y, c.y, d.y, e.y, f.y, g.y, > + h.y, i.y, j.y, k.y, l.y, m.y, n.y, > + clip_max); > + > + > + > + > + dst[yo*dst_pitch+xo] = final; > +} > + > ... > + > +} /* extern "C" */ > -- > 2.25.1 > Best regards, Thomas _______________________________________________ ffmpeg-devel mailing list ffmpeg-devel@ffmpeg.org https://ffmpeg.org/mailman/listinfo/ffmpeg-devel To unsubscribe, visit link above, or email ffmpeg-devel-requ...@ffmpeg.org with subject "unsubscribe".