diff --git a/libavfilter/opencl/nlmeans.cl b/libavfilter/opencl/nlmeans.cl
index 72bd681fd6..6d78a41e46 100644
--- a/libavfilter/opencl/nlmeans.cl
+++ b/libavfilter/opencl/nlmeans.cl
@@ -20,7 +20,7 @@ const sampler_t sampler = (CLK_NORMALIZED_COORDS_FALSE |
                            CLK_ADDRESS_CLAMP_TO_EDGE   |
                            CLK_FILTER_NEAREST);
 
-kernel void horiz_sum(__global uint4 *integral_img,
+kernel void horiz_sum(__global float4 *integral_img,
                       __read_only image2d_t src,
                       int width,
                       int height,
@@ -31,7 +31,7 @@ kernel void horiz_sum(__global uint4 *integral_img,
     int y = get_global_id(0);
     int work_size = get_global_size(0);
 
-    uint4 sum = (uint4)(0);
+    float4 sum = 0.0;
     float4 s2;
     for (int i = 0; i < width; i++) {
         float s1 = read_imagef(src, sampler, (int2)(i, y)).x;
@@ -39,28 +39,25 @@ kernel void horiz_sum(__global uint4 *integral_img,
         s2.y = read_imagef(src, sampler, (int2)(i + dx.y, y + dy.y)).x;
         s2.z = read_imagef(src, sampler, (int2)(i + dx.z, y + dy.z)).x;
         s2.w = read_imagef(src, sampler, (int2)(i + dx.w, y + dy.w)).x;
-        sum += convert_uint4((s1 - s2) * (s1 - s2) * 255 * 255);
+        sum += (s1 - s2) * (s1 - s2);
         integral_img[y * width + i] = sum;
     }
 }
 
-kernel void vert_sum(__global uint4 *integral_img,
-                     __global int *overflow,
+kernel void vert_sum(__global float4 *integral_img,
                      int width,
                      int height)
 {
     int x = get_global_id(0);
-    uint4 sum = 0;
+    float4 sum = 0;
     for (int i = 0; i < height; i++) {
-        if (any((uint4)UINT_MAX - integral_img[i * width + x] < sum))
-            atomic_inc(overflow);
         integral_img[i * width + x] += sum;
         sum = integral_img[i * width + x];
     }
 }
 
 kernel void weight_accum(global float *sum, global float *weight,
-                         global uint4 *integral_img, __read_only image2d_t src,
+                         global float4 *integral_img, __read_only image2d_t src,
                          int width, int height, int p, float h,
                          int4 dx, int4 dy)
 {
@@ -75,16 +72,16 @@ kernel void weight_accum(global float *sum, global float *weight,
     int y = get_global_id(1);
     int4 xoff = x + dx;
     int4 yoff = y + dy;
-    uint4 a = 0, b = 0, c = 0, d = 0;
-    uint4 src_pix = 0;
+    float4 a = 0, b = 0, c = 0, d = 0;
+    float4 src_pix = 0;
 
     // out-of-bounding-box?
     int oobb = (x - p) < 0 || (y - p) < 0 || (y + p) >= height || (x + p) >= width;
 
-    src_pix.x = (int)(255 * read_imagef(src, sampler, (int2)(xoff.x, yoff.x)).x);
-    src_pix.y = (int)(255 * read_imagef(src, sampler, (int2)(xoff.y, yoff.y)).x);
-    src_pix.z = (int)(255 * read_imagef(src, sampler, (int2)(xoff.z, yoff.z)).x);
-    src_pix.w = (int)(255 * read_imagef(src, sampler, (int2)(xoff.w, yoff.w)).x);
+    src_pix.x = read_imagef(src, sampler, (int2)(xoff.x, yoff.x)).x;
+    src_pix.y = read_imagef(src, sampler, (int2)(xoff.y, yoff.y)).x;
+    src_pix.z = read_imagef(src, sampler, (int2)(xoff.z, yoff.z)).x;
+    src_pix.w = read_imagef(src, sampler, (int2)(xoff.w, yoff.w)).x;
     if (!oobb) {
         a = integral_img[(y - p) * width + x - p];
         b = integral_img[(y + p) * width + x - p];
@@ -93,7 +90,7 @@ kernel void weight_accum(global float *sum, global float *weight,
     }
 
     float4 patch_diff = convert_float4(d + a - c - b);
-    float4 w = native_exp(-patch_diff / (h * h));
+    float4 w = native_exp(-patch_diff * (float4)(255.0f * 255.0f) / (h * h));
     float w_sum = w.x + w.y + w.z + w.w;
     weight[y * width + x] += w_sum;
     sum[y * width + x] += dot(w, convert_float4(src_pix));
@@ -109,7 +106,7 @@ kernel void average(__write_only image2d_t dst,
     float w = weight[y * dim.x + x];
     float s = sum[y * dim.x + x];
     float src_pix = read_imagef(src, sampler, (int2)(x, y)).x;
-    float r = (s + src_pix * 255) / (1.0f + w) / 255.0f;
+    float r = (s + src_pix) / (1.0f + w);
     if (x < dim.x && y < dim.y)
         write_imagef(dst, (int2)(x, y), (float4)(r, 0.0f, 0.0f, 1.0f));
 }
diff --git a/libavfilter/vf_nlmeans_opencl.c b/libavfilter/vf_nlmeans_opencl.c
index e57b5e0873..0b69f3b6c4 100644
--- a/libavfilter/vf_nlmeans_opencl.c
+++ b/libavfilter/vf_nlmeans_opencl.c
@@ -30,11 +30,9 @@
 #include "opencl_source.h"
 #include "video.h"
 
-// TODO:
-//      the integral image may overflow 32bit, consider using 64bit
-
 static const enum AVPixelFormat supported_formats[] = {
     AV_PIX_FMT_YUV420P,
+    AV_PIX_FMT_YUV420P16LE,
     AV_PIX_FMT_YUV444P,
     AV_PIX_FMT_GBRP,
 };
@@ -59,7 +57,6 @@ typedef struct NLMeansOpenCLContext {
     cl_mem                integral_img;
     cl_mem                weight;
     cl_mem                sum;
-    cl_mem                overflow; // overflow in integral image?
     double                sigma;
     float                 h;
     int                   chroma_w;
@@ -129,7 +126,7 @@ static int nlmeans_opencl_init(AVFilterContext *avctx, int width, int height)
                      "average kernel %d.\n", cle);
 
     ctx->integral_img = clCreateBuffer(ctx->ocf.hwctx->context, 0,
-                                       4 * width * height * sizeof(cl_int),
+                                       4 * width * height * sizeof(cl_float),
                                        NULL, &cle);
     CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create "
                      "integral image %d.\n", cle);
@@ -144,11 +141,6 @@ static int nlmeans_opencl_init(AVFilterContext *avctx, int width, int height)
     CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create "
                      "sum buffer %d.\n", cle);
 
-    ctx->overflow = clCreateBuffer(ctx->ocf.hwctx->context, 0,
-                                   sizeof(cl_int), NULL, &cle);
-    CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create "
-                     "overflow buffer %d.\n", cle);
-
     ctx->initialised = 1;
     return 0;
 
@@ -161,7 +153,6 @@ fail:
     CL_RELEASE_MEMORY(ctx->integral_img);
     CL_RELEASE_MEMORY(ctx->weight);
     CL_RELEASE_MEMORY(ctx->sum);
-    CL_RELEASE_MEMORY(ctx->overflow);
 
     CL_RELEASE_QUEUE(ctx->command_queue);
     return err;
@@ -239,9 +230,8 @@ static int nlmeans_plane(AVFilterContext *avctx, cl_mem dst, cl_mem src,
         // vertical pass
         // integral(x, y) = sum(integral(x, v)) for v in [0, y]
         CL_SET_KERNEL_ARG(ctx->vert_kernel, 0, cl_mem, &ctx->integral_img);
-        CL_SET_KERNEL_ARG(ctx->vert_kernel, 1, cl_mem, &ctx->overflow);
-        CL_SET_KERNEL_ARG(ctx->vert_kernel, 2, cl_int, &width);
-        CL_SET_KERNEL_ARG(ctx->vert_kernel, 3, cl_int, &height);
+        CL_SET_KERNEL_ARG(ctx->vert_kernel, 1, cl_int, &width);
+        CL_SET_KERNEL_ARG(ctx->vert_kernel, 2, cl_int, &height);
         cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->vert_kernel,
                                      1, NULL, worksize2, NULL, 0, NULL, NULL);
         CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue vert_kernel: %d.\n",
@@ -293,8 +283,7 @@ static int nlmeans_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input)
     const AVPixFmtDescriptor *desc;
     enum AVPixelFormat in_format;
     cl_mem src, dst;
-    const cl_int zero = 0;
-    int w, h, err, cle, overflow, p, patch, research;
+    int w, h, err, cle, p, patch, research;
 
     av_log(ctx, AV_LOG_DEBUG, "Filter input: %s, %ux%u (%"PRId64").\n",
            av_get_pix_fmt_name(input->format),
@@ -331,11 +320,6 @@ static int nlmeans_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input)
             goto fail;
     }
 
-    cle = clEnqueueWriteBuffer(ctx->command_queue, ctx->overflow, CL_FALSE,
-                               0, sizeof(cl_int), &zero, 0, NULL, NULL);
-    CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to initialize overflow"
-                     "detection buffer %d.\n", cle);
-
     for (p = 0; p < FF_ARRAY_ELEMS(output->data); p++) {
         src = (cl_mem) input->data[p];
         dst = (cl_mem) output->data[p];
@@ -351,17 +335,10 @@ static int nlmeans_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input)
         if (err < 0)
             goto fail;
     }
-    // overflow occurred?
-    cle = clEnqueueReadBuffer(ctx->command_queue, ctx->overflow, CL_FALSE,
-                              0, sizeof(cl_int), &overflow, 0, NULL, NULL);
-    CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to read overflow: %d.\n", cle);
 
     cle = clFinish(ctx->command_queue);
     CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to finish kernel: %d.\n", cle);
 
-    if (overflow > 0)
-        av_log(avctx, AV_LOG_ERROR, "integral image overflow %d\n", overflow);
-
     av_frame_free(&input);
 
     av_log(ctx, AV_LOG_DEBUG, "Filter output: %s, %ux%u (%"PRId64").\n",
@@ -390,7 +367,6 @@ static av_cold void nlmeans_opencl_uninit(AVFilterContext *avctx)
     CL_RELEASE_MEMORY(ctx->integral_img);
     CL_RELEASE_MEMORY(ctx->weight);
     CL_RELEASE_MEMORY(ctx->sum);
-    CL_RELEASE_MEMORY(ctx->overflow);
 
     CL_RELEASE_QUEUE(ctx->command_queue);
 
