             return res;
         }
     };
-    template <> struct SafeDiv<4>
-    {
-        template <typename T>
+    template <typename T> struct SafeDiv<T, 4>
+    {
         __device__ __forceinline__ static T op(const T& a, const T& b)
         {
             T res;


        __device__ __forceinline__ DstType operator ()(SrcType a) const
        {
            return saturate_cast<DstType>(SafeDiv<VecTraits<ScalarType>::cn>::op(saturate_cast<ScalarType>(a), val));
        }
    };



namespace
{
    template <typename T, int cn> struct SafeDiv;
    template <typename T> struct SafeDiv<T, 1>
    {
        __device__ __forceinline__ static T op(T a, T b)
        {
            return b != 0 ? a / b : 0;
        }
    };
    template <typename T> struct SafeDiv<T, 2>
    {
        __device__ __forceinline__ static T op(const T& a, const T& b)
        {
            T res;
