diff options
author | Philip Langdale <philipl@overt.org> | 2019-02-20 19:57:52 -0800 |
---|---|---|
committer | Timo Rothenpieler <timo@rothenpieler.org> | 2019-02-27 18:03:17 +0100 |
commit | b4c9c09915de8ffaa4a2d2606e85729afa6c4e8e (patch) | |
tree | cb701f8d84c90c8f47f10045df00620031401d73 /libavfilter/vf_thumbnail_cuda.cu | |
parent | 2544c7ea67ca9521c5de36396bc9ac7058223742 (diff) | |
download | ffmpeg-b4c9c09915de8ffaa4a2d2606e85729afa6c4e8e.tar.gz |
avfilter/vf_thumbnail_cuda: Switch to using ffnvcodec
This change switches the vf_thumbnail_cuda filter from using the
full cuda sdk to using the ffnvcodec headers and loader.
Most of the change is a direct mapping, but I also switched from
using texture references to using texture objects. This is supposed
to be the preferred way of using textures, and the texture object API
is the one I added to ffnvcodec.
Signed-off-by: Philip Langdale <philipl@overt.org>
Signed-off-by: Timo Rothenpieler <timo@rothenpieler.org>
Diffstat (limited to 'libavfilter/vf_thumbnail_cuda.cu')
-rw-r--r-- | libavfilter/vf_thumbnail_cuda.cu | 25 |
1 files changed, 12 insertions, 13 deletions
diff --git a/libavfilter/vf_thumbnail_cuda.cu b/libavfilter/vf_thumbnail_cuda.cu index 98fad4303a..c73e49fbc6 100644 --- a/libavfilter/vf_thumbnail_cuda.cu +++ b/libavfilter/vf_thumbnail_cuda.cu @@ -22,55 +22,54 @@ 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); } |