Skip to content

[cudaResize] Added some interpolation filters. linear, area, cubic, lanczos4, spline36. #37

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
wants to merge 7 commits into
base: master
Choose a base branch
from
Open

[cudaResize] Added some interpolation filters. linear, area, cubic, lanczos4, spline36. #37

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
7 commits
Select commit Hold shift + click to select a range
c8ea002
[cudaResize] add interporation filter types. 'linear' & 'area'.
TE-0000123718 Sep 1, 2020
02f2192
[examples] Change sample code of resize function in cuda-examples.py.
TE-0000123718 Sep 4, 2020
31dc9bc
Use 'cast_vec' instead of 'from_float4'. remove 'from_float4'.
HiroshiYAMA Sep 5, 2020
c12b135
[cudaResize] Change adaptive selection logic for 'AREA' filter.
HiroshiYAMA Sep 5, 2020
346d90e
[cudaResize] Change for a little performance improvement.
HiroshiYAMA Sep 6, 2020
a1a29b6
[cudaResize] Add interpolation filter types. cubic, lanczos4, spline36.
HiroshiYAMA Sep 9, 2020
e903773
Update comment message.
HiroshiYAMA Sep 10, 2020
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
104 changes: 104 additions & 0 deletions cuda/cudaMath.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -24,6 +24,9 @@
#define __CUDA_HELPER_MATH_H_

#include "cuda_runtime.h"
#include <limits>
#include <cfloat>
#include <cmath>

////////////////////////////////////////////////////////////////////////////////
/// @name Vector Math
Expand Down Expand Up @@ -125,6 +128,10 @@ inline __host__ __device__ uint2 make_uint2(int2 a)
return make_uint2(uint(a.x), uint(a.y));
}

inline __host__ __device__ float3 make_float3(uchar s)
{
return make_float3(s, s, s);
}
inline __host__ __device__ float3 make_float3(float s)
{
return make_float3(s, s, s);
Expand Down Expand Up @@ -245,6 +252,10 @@ inline __host__ __device__ uchar3 make_uchar3(float4 a)
return make_uchar3(a.x, a.y, a.z);
}

inline __host__ __device__ float4 make_float4(uchar s)
{
return make_float4(s, s, s, s);
}
inline __host__ __device__ float4 make_float4(float s)
{
return make_float4(s, s, s, s);
Expand Down Expand Up @@ -520,6 +531,27 @@ inline __host__ __device__ void operator+=(uint3 &a, uint b)
a.z += b;
}

inline __host__ __device__ uchar3 operator+(uchar3 a, uchar3 b)
{
return make_uchar3(a.x + b.x, a.y + b.y, a.z + b.z);
}
inline __host__ __device__ void operator+=(uchar3 &a, uchar3 b)
{
a.x += b.x;
a.y += b.y;
a.z += b.z;
}
inline __host__ __device__ uchar3 operator+(uchar3 a, uchar b)
{
return make_uchar3(a.x + b, a.y + b, a.z + b);
}
inline __host__ __device__ void operator+=(uchar3 &a, uchar b)
{
a.x += b;
a.y += b;
a.z += b;
}

inline __host__ __device__ int3 operator+(int b, int3 a)
{
return make_int3(a.x + b, a.y + b, a.z + b);
Expand All @@ -532,6 +564,10 @@ inline __host__ __device__ float3 operator+(float b, float3 a)
{
return make_float3(a.x + b, a.y + b, a.z + b);
}
inline __host__ __device__ uchar3 operator+(uchar b, uchar3 a)
{
return make_uchar3(a.x + b, a.y + b, a.z + b);
}

inline __host__ __device__ float4 operator+(float4 a, float4 b)
{
Expand Down Expand Up @@ -614,6 +650,33 @@ inline __host__ __device__ void operator+=(uint4 &a, uint b)
a.w += b;
}

inline __host__ __device__ uchar4 operator+(uchar4 a, uchar4 b)
{
return make_uchar4(a.x + b.x, a.y + b.y, a.z + b.z, a.w + b.w);
}
inline __host__ __device__ void operator+=(uchar4 &a, uchar4 b)
{
a.x += b.x;
a.y += b.y;
a.z += b.z;
a.w += b.w;
}
inline __host__ __device__ uchar4 operator+(uchar4 a, uchar b)
{
return make_uchar4(a.x + b, a.y + b, a.z + b, a.w + b);
}
inline __host__ __device__ uchar4 operator+(uchar b, uchar4 a)
{
return make_uchar4(a.x + b, a.y + b, a.z + b, a.w + b);
}
inline __host__ __device__ void operator+=(uchar4 &a, uchar b)
{
a.x += b;
a.y += b;
a.z += b;
a.w += b;
}

////////////////////////////////////////////////////////////////////////////////
// subtract
////////////////////////////////////////////////////////////////////////////////
Expand Down Expand Up @@ -986,6 +1049,20 @@ inline __host__ __device__ void operator*=(uint3 &a, uint b)
a.y *= b;
a.z *= b;
}
inline __host__ __device__ uchar3 operator*(uchar3 a, float b)
{
return make_uchar3(a.x * b, a.y * b, a.z * b);
}
inline __host__ __device__ uchar3 operator*(float b, uchar3 a)
{
return make_uchar3(b * a.x, b * a.y, b * a.z);
}
inline __host__ __device__ void operator*=(uchar3 &a, float b)
{
a.x *= b;
a.y *= b;
a.z *= b;
}

inline __host__ __device__ float4 operator*(float4 a, float4 b)
{
Expand Down Expand Up @@ -1068,6 +1145,33 @@ inline __host__ __device__ void operator*=(uint4 &a, uint b)
a.w *= b;
}

inline __host__ __device__ uchar4 operator*(uchar4 a, uchar4 b)
{
return make_uchar4(a.x * b.x, a.y * b.y, a.z * b.z, a.w * b.w);
}
inline __host__ __device__ void operator*=(uchar4 &a, uchar4 b)
{
a.x *= b.x;
a.y *= b.y;
a.z *= b.z;
a.w *= b.w;
}
inline __host__ __device__ uchar4 operator*(uchar4 a, uchar b)
{
return make_uchar4(a.x * b, a.y * b, a.z * b, a.w * b);
}
inline __host__ __device__ uchar4 operator*(uchar b, uchar4 a)
{
return make_uchar4(b * a.x, b * a.y, b * a.z, b * a.w);
}
inline __host__ __device__ void operator*=(uchar4 &a, uchar b)
{
a.x *= b;
a.y *= b;
a.z *= b;
a.w *= b;
}

////////////////////////////////////////////////////////////////////////////////
// divide
////////////////////////////////////////////////////////////////////////////////
Expand Down
Loading