sse_cuda.h

Go to the documentation of this file.

#include "simd_functions.h"

#if defined( __CUDACC__ )

typedef uint4 m128i;

__device__ m128i mm_cached_load_si128 (m128i *p)
{
    return __ldg(p);
}
__device__ m128i mm_load_si128 (m128i *p)
{
    return *p;
}
__device__ void mm_store_si128 (m128i *p, m128i a)
{
    *p = a;
}
__device__ m128i mm_set_epi8 (char b15, char b14, char b13, char b12,char b11, char b10, char b9, char b8, char b7, char b6, char b5, char b4, char b3, char b2, char b1, char b0)
{
    return make_uint4(
        unsigned int(unsigned char(b0)+128)       |
        unsigned int(unsigned char(b1)+128) << 8  |
        unsigned int(unsigned char(b2)+128) << 16 |
        unsigned int(unsigned char(b3)+128) << 24,

        unsigned int(unsigned char(b4)+128)       |
        unsigned int(unsigned char(b5)+128) << 8  |
        unsigned int(unsigned char(b6)+128) << 16 |
        unsigned int(unsigned char(b7)+128) << 24,

        unsigned int(unsigned char(b8)+128)        |
        unsigned int(unsigned char(b9)+128)  << 8  |
        unsigned int(unsigned char(b10)+128) << 16 |
        unsigned int(unsigned char(b11)+128) << 24,

        unsigned int(unsigned char(b12)+128)       |
        unsigned int(unsigned char(b13)+128) << 8  |
        unsigned int(unsigned char(b14)+128) << 16 |
        unsigned int(unsigned char(b15)+128) << 24 );
}
__device__ m128i mm_set1_epi32 (int i)
{
    return make_uint4( i, i, i, i );
}
__device__ m128i mm_set1_epi8 (char b)
{
    const unsigned int ub = unsigned int(unsigned char(b)+128);
    const unsigned int uw = ub | ub << 8 | ub << 16 | ub << 24;
    return make_uint4( uw, uw, uw, uw );
}
__device__ m128i mm_adds_epu8 (m128i a, m128i b)
{
    return make_uint4(
        vaddus4( a.x, b.x ),
        vaddus4( a.y, b.y ),
        vaddus4( a.z, b.z ),
        vaddus4( a.w, b.w ) );
}
__device__ m128i mm_subs_epu8 (m128i a, m128i b)
{
    return make_uint4(
        vsubus4( a.x, b.x ),
        vsubus4( a.y, b.y ),
        vsubus4( a.z, b.z ),
        vsubus4( a.w, b.w ) );
}

__device__ m128i mm_max_epu8 (m128i a, m128i b)
{
    return make_uint4(
        vmaxu4( a.x, b.x ),
        vmaxu4( a.y, b.y ),
        vmaxu4( a.z, b.z ),
        vmaxu4( a.w, b.w ) );
}
__device__ m128i mm_min_epu8 (m128i a, m128i b)
{
    return make_uint4(
        vminu4( a.x, b.x ),
        vminu4( a.y, b.y ),
        vminu4( a.z, b.z ),
        vminu4( a.w, b.w ) );
}

__device__ m128i mm_cmpeq_epi8 (m128i a, m128i b)
{
    return make_uint4(
        vcmpeq4( a.x, b.x ),
        vcmpeq4( a.y, b.y ),
        vcmpeq4( a.z, b.z ),
        vcmpeq4( a.w, b.w ) );
}

__device__ int mm_movemask_epi8 (m128i a)
{
    return  ((a.x & (1u << 31)) >> 16) |
            ((a.x & (1u << 23)) >>  9) |
            ((a.x & (1u << 15)) >>  2) |
            ((a.x & (1u <<  7)) <<  5) |
            ((a.y & (1u << 31)) >> 20) |
            ((a.y & (1u << 23)) >> 13) |
            ((a.y & (1u << 15)) >>  6) |
            ((a.y & (1u <<  7)) <<  1) |
            ((a.z & (1u << 31)) >> 24) |
            ((a.z & (1u << 23)) >> 17) |
            ((a.z & (1u << 15)) >> 10) |
            ((a.z & (1u <<  7)) >>  3) |
            ((a.w & (1u << 31)) >> 28) |
            ((a.w & (1u << 23)) >> 21) |
            ((a.w & (1u << 15)) >> 14) |
            ((a.w & (1u <<  7)) >>  7);
}

template <unsigned int imm>
struct mm_slli_si128_dispatcher {};

template <>
struct mm_slli_si128_dispatcher<1>
{
    __device__ static m128i dispatch(const m128i a)
    {
        return make_uint4(
            (a.x << 8),
            (a.y << 8) | (a.x >> 24),
            (a.z << 8) | (a.y >> 24),
            (a.w << 8) | (a.z >> 24) );
    }
};

template <>
struct mm_slli_si128_dispatcher<2>
{
    __device__ static m128i dispatch(const m128i a)
    {
        return make_uint4(
            (a.x << 16),
            (a.y << 16) | (a.x >> 16),
            (a.z << 16) | (a.y >> 16),
            (a.w << 16) | (a.z >> 16) );
    }
};

template <>
struct mm_slli_si128_dispatcher<3>
{
    __device__ static m128i dispatch(const m128i a)
    {
        return make_uint4(
            (a.x << 24),
            (a.y << 24) | (a.x >> 8),
            (a.z << 24) | (a.y >> 8),
            (a.w << 24) | (a.z >> 8) );
    }
};

template <>
struct mm_slli_si128_dispatcher<4>
{
    __device__ static m128i dispatch(const m128i a)
    {
        return make_uint4(
            0,
            a.x,
            a.y,
            a.z );
    }
};

template <>
struct mm_slli_si128_dispatcher<8>
{
    __device__ static m128i dispatch(const m128i a)
    {
        return make_uint4(
            0,
            0,
            a.x,
            a.y );
    }
};

template <>
struct mm_slli_si128_dispatcher<12>
{
    __device__ static m128i dispatch(const m128i a)
    {
        return make_uint4(
            0,
            0,
            0,
            a.x );
    }
};

#define mm_slli_si128( a, imm ) mm_slli_si128_dispatcher<imm>::dispatch( a )

template <unsigned int imm>
struct mm_srli_si128_dispatcher {};

template <>
struct mm_srli_si128_dispatcher<1>
{
    __device__ static m128i dispatch(const m128i a)
    {
        return make_uint4(
            (a.x >> 8) | (a.y & 255u),
            (a.y >> 8) | (a.z & 255u),
            (a.z >> 8) | (a.w & 255u),
            (a.w >> 8) );
    }
};

template <>
struct mm_srli_si128_dispatcher<2>
{
    __device__ static m128i dispatch(const m128i a)
    {
        return make_uint4(
            (a.x >> 16) | (a.y & 0x0000FFFFu),
            (a.y >> 16) | (a.z & 0x0000FFFFu),
            (a.z >> 16) | (a.w & 0x0000FFFFu),
            (a.w >> 16) );
    }
};

template <>
struct mm_srli_si128_dispatcher<3>
{
    __device__ static m128i dispatch(const m128i a)
    {
        return make_uint4(
            (a.x >> 24) | (a.y & 0x00FFFFFFu),
            (a.y >> 24) | (a.z & 0x00FFFFFFu),
            (a.z >> 24) | (a.w & 0x00FFFFFFu),
            (a.w >> 24) );
    }
};

template <>
struct mm_srli_si128_dispatcher<4>
{
    __device__ static m128i dispatch(const m128i a)
    {
        return make_uint4(
            a.y,
            a.z,
            a.w,
            0 );
    }
};

template <>
struct mm_srli_si128_dispatcher<8>
{
    __device__ static m128i dispatch(const m128i a)
    {
        return make_uint4(
            a.z,
            a.w,
            0,
            0 );
    }
};

template <>
struct mm_srli_si128_dispatcher<12>
{
    __device__ static m128i dispatch(const m128i a)
    {
        return make_uint4(
            a.w,
            0,
            0,
            0 );
    }
};

#define mm_srli_si128( a, imm ) mm_srli_si128_dispatcher<imm>::dispatch( a )

template <unsigned int imm> struct mm_extract_epi16_dispatcher {};
template <> struct mm_extract_epi16_dispatcher<0> { __device__ static int dispatch(const m128i a) { return a.w >> 16; } };
template <> struct mm_extract_epi16_dispatcher<1> { __device__ static int dispatch(const m128i a) { return a.w & 0x0000FFFF; } };
template <> struct mm_extract_epi16_dispatcher<2> { __device__ static int dispatch(const m128i a) { return a.z >> 16; } };
template <> struct mm_extract_epi16_dispatcher<3> { __device__ static int dispatch(const m128i a) { return a.z & 0x0000FFFF; } };
template <> struct mm_extract_epi16_dispatcher<4> { __device__ static int dispatch(const m128i a) { return a.y >> 16; } };
template <> struct mm_extract_epi16_dispatcher<5> { __device__ static int dispatch(const m128i a) { return a.y & 0x0000FFFF; } };
template <> struct mm_extract_epi16_dispatcher<6> { __device__ static int dispatch(const m128i a) { return a.x >> 16; } };
template <> struct mm_extract_epi16_dispatcher<7> { __device__ static int dispatch(const m128i a) { return a.x & 0x0000FFFF; } };

#define mm_extract_epi16( a, imm ) mm_extract_epi16_dispatcher<imm>::dispatch( a )

__global__ void sse_test_kernel()
{
    m128i zero, a, b;

    zero = mm_set1_epi32(0);
    a    = mm_set1_epi8(0xFF);
    b    = mm_cmpeq_epi8( zero, zero );
    if (a.x != b.x ||
        a.y != b.y ||
        a.z != b.z ||
        a.w != b.w) printf("error\n");

    b = mm_set_epi8( 15, 14, 13, 12, 11, 10, 9, 8, 7, 6, 5, 4, 3, 2, 1, 0 );

    int h0 = mm_extract_epi16( b, 0 );
}

#endif // __CUDACC__