X-Git-Url: https://git.sesse.net/?a=blobdiff_plain;f=libavfilter%2Fvf_thumbnail_cuda.cu;h=d4d4d791f6819eb7743c69e22529b14e3cccb90b;hb=3c77584be8fa6833b5cbf9b064ef1120a54fa61f;hp=98fad4303a74a510d38bae400dd300dfebcb477f;hpb=8e789d244cc946bc350672eeb02453918b21a09f;p=ffmpeg

diff --git a/libavfilter/vf_thumbnail_cuda.cu b/libavfilter/vf_thumbnail_cuda.cu
index 98fad4303a7..d4d4d791f68 100644
--- a/libavfilter/vf_thumbnail_cuda.cu
+++ b/libavfilter/vf_thumbnail_cuda.cu
@@ -22,57 +22,56 @@
 
 extern "C" {
 
-texture<unsigned char, 2> uchar_tex;
-texture<uchar2, 2>  uchar2_tex;
-texture<unsigned short, 2> ushort_tex;
-texture<ushort2, 2>  ushort2_tex;
-
-__global__ void Thumbnail_uchar(int *histogram, int src_width, int src_height)
+__global__ void Thumbnail_uchar(cudaTextureObject_t uchar_tex,
+                                int *histogram, int src_width, int src_height)
 {
     int x = blockIdx.x * blockDim.x + threadIdx.x;
     int y = blockIdx.y * blockDim.y + threadIdx.y;
     if (y < src_height && x < src_width)
     {
-        unsigned char pixel = tex2D(uchar_tex, x, y);
+        unsigned char pixel = tex2D<unsigned char>(uchar_tex, x, y);
         atomicAdd(&histogram[pixel], 1);
     }
 }
 
-__global__ void Thumbnail_uchar2(int *histogram, int src_width, int src_height)
+__global__ void Thumbnail_uchar2(cudaTextureObject_t uchar2_tex,
+                                 int *histogram, int src_width, int src_height)
 {
     int x = blockIdx.x * blockDim.x + threadIdx.x;
     int y = blockIdx.y * blockDim.y + threadIdx.y;
 
     if (y < src_height && x < src_width)
     {
-        uchar2 pixel = tex2D(uchar2_tex, x, y);
+        uchar2 pixel = tex2D<uchar2>(uchar2_tex, x, y);
         atomicAdd(&histogram[pixel.x], 1);
         atomicAdd(&histogram[256 + pixel.y], 1);
     }
 }
 
-__global__ void Thumbnail_ushort(int *histogram, int src_width, int src_height)
+__global__ void Thumbnail_ushort(cudaTextureObject_t ushort_tex,
+                                 int *histogram, int src_width, int src_height)
 {
     int x = blockIdx.x * blockDim.x + threadIdx.x;
     int y = blockIdx.y * blockDim.y + threadIdx.y;
 
     if (y < src_height && x < src_width)
     {
-        unsigned short pixel = (tex2D(ushort_tex, x, y) + 128) >> 8;
+        unsigned short pixel = (tex2D<unsigned short>(ushort_tex, x, y) + 128) >> 8;
         atomicAdd(&histogram[pixel], 1);
     }
 }
 
-__global__ void Thumbnail_ushort2(int *histogram, int src_width, int src_height)
+__global__ void Thumbnail_ushort2(cudaTextureObject_t ushort2_tex,
+                                  int *histogram, int src_width, int src_height)
 {
     int x = blockIdx.x * blockDim.x + threadIdx.x;
     int y = blockIdx.y * blockDim.y + threadIdx.y;
 
     if (y < src_height && x < src_width)
     {
-        ushort2 pixel = tex2D(ushort2_tex, x, y);
+        ushort2 pixel = tex2D<ushort2>(ushort2_tex, x, y);
         atomicAdd(&histogram[(pixel.x + 128) >> 8], 1);
-        atomicAdd(&histogram[256 + (pixel.y + 128) >> 8], 1);
+        atomicAdd(&histogram[256 + ((pixel.y + 128) >> 8)], 1);
     }
 }