Bit-manipulation intrinsics#

Defined in code/simulator.h.

This page documents 48 symbols: 1 macros, 47 functions.

Macros#

Signature

Description

Availability

Location

Example

SHARED_EXTERN(type, name)

Expands to type* name = static_cast<type*>(::simDynamicShared())

sim + nvcc

simulator.h:949

View

Functions#

Signature

Description

Availability

Location

Example

inline unsigned __brev(unsigned x)

sim only

simulator.h:808

inline unsigned long long __brevll(unsigned long long x)

sim only

simulator.h:815

inline unsigned __byte_perm(unsigned x, unsigned y, unsigned sel)

Assemble 4 bytes selected by the nibbles of sel out of the 8-byte pair {x,y}.

sim only

simulator.h:820

inline int __clz (int x)

sim only

simulator.h:803

inline int __clzll (long long x)

sim only

simulator.h:804

inline int __dp2a_hi(int a, int b, int c)

sim only

simulator.h:882

inline unsigned __dp2a_hi(unsigned a, unsigned b, unsigned c)

sim only

simulator.h:890

inline int __dp2a_lo(int a, int b, int c)

__dp2a_lo uses the LOW halves of a’s shorts, _hi the high halves.

sim only

simulator.h:877

inline unsigned __dp2a_lo(unsigned a, unsigned b, unsigned c)

sim only

simulator.h:887

inline int __dp4a(int a, int b, int c)

sm_61 4-way byte / 2-way short dot products with accumulate.

sim only

simulator.h:855

inline unsigned __dp4a(unsigned a, unsigned b, unsigned c)

sim only

simulator.h:862

inline int __dp4a(char4 a, char4 b, int c)

sim only

simulator.h:869

inline unsigned __dp4a(uchar4 a, uchar4 b, unsigned c)

sim only

simulator.h:872

inline int __ffs (int x)

__ffs is 1-based and returns 0 for an all-zero word (CUDA semantics).

sim only

simulator.h:806

inline int __ffsll (long long x)

sim only

simulator.h:807

inline int __float_as_int(float x)

sim only

simulator.h:897

inline unsigned __funnelshift_l(unsigned lo, unsigned hi, unsigned shift)

Concatenate {hi:lo} to 64 bits, shift, take the requested 32-bit window.

sim only

simulator.h:836

inline unsigned __funnelshift_lc(unsigned lo, unsigned hi, unsigned shift)

sim only

simulator.h:840

inline unsigned __funnelshift_r(unsigned lo, unsigned hi, unsigned shift)

sim only

simulator.h:845

inline unsigned __funnelshift_rc(unsigned lo, unsigned hi, unsigned shift)

sim only

simulator.h:849

inline float __int_as_float(int x)

sim only

simulator.h:896

inline int __popc (unsigned x)

sim only

simulator.h:801

inline int __popcll(unsigned long long x)

sim only

simulator.h:802

inline int __sad(int x, int y, int z)

sim only

simulator.h:893

inline unsigned __usad(unsigned x, unsigned y, unsigned z)

sim only

simulator.h:894

inline long long llmax(long long a, long long b)

sim only

simulator.h:919

inline long long llmin(long long a, long long b)

sim only

simulator.h:918

inline int max(int a, int b)

Device math provides max/min overloads that host <algorithm> only exposes as std::max/std::min.

sim only

simulator.h:903

View

inline unsigned max(unsigned a, unsigned b)

sim only

simulator.h:904

View

inline long long max(long long a, long long b)

sim only

simulator.h:905

View

inline unsigned long long max(unsigned long long a, unsigned long long b)

sim only

simulator.h:906

View

inline float max(float a, float b)

sim only

simulator.h:907

View

inline double max(double a, double b)

sim only

simulator.h:908

View

inline std::common_type_t<A, B> max(A a, B b)

sim only

simulator.h:931

View

inline int min(int a, int b)

sim only

simulator.h:909

View

inline unsigned min(unsigned a, unsigned b)

sim only

simulator.h:910

View

inline long long min(long long a, long long b)

sim only

simulator.h:911

View

inline unsigned long long min(unsigned long long a, unsigned long long b)

sim only

simulator.h:912

View

inline float min(float a, float b)

sim only

simulator.h:913

View

inline double min(double a, double b)

sim only

simulator.h:914

View

inline std::common_type_t<A, B> min(A a, B b)

sim only

simulator.h:938

View

inline void sim_sync()

__syncthreads(): park at the block-wide barrier; the scheduler releases it only once every live thread in the block has arrived.

sim only

simulator.h:953

inline void* simDynamicShared()

Dynamic (runtime-sized) shared memory.

sim only

simulator.h:948

inline unsigned long long ullmax(unsigned long long a, unsigned long long b)

sim only

simulator.h:921

inline unsigned long long ullmin(unsigned long long a, unsigned long long b)

sim only

simulator.h:920

inline unsigned umax (unsigned a, unsigned b)

sim only

simulator.h:917

inline unsigned umin (unsigned a, unsigned b)

CUDA also exposes the explicitly-typed integer min/max intrinsics by name.

sim only

simulator.h:916

View

Examples#

Code from the repository’s example corpus (code\*.cu); the highlighted line is the call site. These blocks are what the View links above open.

SHARED_EXTERN (test_shared.cu:52)
42    __syncthreads();
43    for (unsigned st = blockDim.x / 2; st > 0; st >>= 1) {
44        if (tid < st) s[tid] += s[tid + st];
45        __syncthreads();
46    }
47    if (tid == 0) outBlk[blockIdx.x] = s[0];
48}
49
50// Same reduction using DYNAMIC shared memory (size passed at launch).
51__global__ void kDynReduce(const float* in, float* outBlk) {
52    SHARED_EXTERN(float, s);
53    const unsigned tid = threadIdx.x;
54    const unsigned i   = blockIdx.x * blockDim.x + tid;
55    s[tid] = in[i];
56    __syncthreads();
57    for (unsigned st = blockDim.x / 2; st > 0; st >>= 1) {
58        if (tid < st) s[tid] += s[tid + st];
59        __syncthreads();
60    }
61    if (tid == 0) outBlk[blockIdx.x] = s[0];
62}
max (test_web_cs_atomics.cu:69)
59    val = 0;
60    for (int i = 0; i < len; ++i) val -= 10;
61    if (val != gpuData[1]) { std::printf("atomicSub failed\n"); return false; }
62
63    bool found = false;
64    for (int i = 0; i < len; ++i) if (i == gpuData[2]) { found = true; break; }
65    if (!found) { std::printf("atomicExch failed\n"); return false; }
66
67    val = -(1 << 8);
68    for (int i = 0; i < len; ++i) val = max(val, i);
69    if (val != gpuData[3]) { std::printf("atomicMax failed\n"); return false; }
70
71    val = 1 << 8;
72    for (int i = 0; i < len; ++i) val = min(val, i);
73    if (val != gpuData[4]) { std::printf("atomicMin failed\n"); return false; }
74
75    int limit = 17;
76    val = 0;
77    for (int i = 0; i < len; ++i) val = (val >= limit) ? 0 : val + 1;
78    if (val != gpuData[5]) { std::printf("atomicInc failed\n"); return false; }
max (test_web_cs_atomics.cu:69)
59    val = 0;
60    for (int i = 0; i < len; ++i) val -= 10;
61    if (val != gpuData[1]) { std::printf("atomicSub failed\n"); return false; }
62
63    bool found = false;
64    for (int i = 0; i < len; ++i) if (i == gpuData[2]) { found = true; break; }
65    if (!found) { std::printf("atomicExch failed\n"); return false; }
66
67    val = -(1 << 8);
68    for (int i = 0; i < len; ++i) val = max(val, i);
69    if (val != gpuData[3]) { std::printf("atomicMax failed\n"); return false; }
70
71    val = 1 << 8;
72    for (int i = 0; i < len; ++i) val = min(val, i);
73    if (val != gpuData[4]) { std::printf("atomicMin failed\n"); return false; }
74
75    int limit = 17;
76    val = 0;
77    for (int i = 0; i < len; ++i) val = (val >= limit) ? 0 : val + 1;
78    if (val != gpuData[5]) { std::printf("atomicInc failed\n"); return false; }
max (test_web_cs_atomics.cu:69)
59    val = 0;
60    for (int i = 0; i < len; ++i) val -= 10;
61    if (val != gpuData[1]) { std::printf("atomicSub failed\n"); return false; }
62
63    bool found = false;
64    for (int i = 0; i < len; ++i) if (i == gpuData[2]) { found = true; break; }
65    if (!found) { std::printf("atomicExch failed\n"); return false; }
66
67    val = -(1 << 8);
68    for (int i = 0; i < len; ++i) val = max(val, i);
69    if (val != gpuData[3]) { std::printf("atomicMax failed\n"); return false; }
70
71    val = 1 << 8;
72    for (int i = 0; i < len; ++i) val = min(val, i);
73    if (val != gpuData[4]) { std::printf("atomicMin failed\n"); return false; }
74
75    int limit = 17;
76    val = 0;
77    for (int i = 0; i < len; ++i) val = (val >= limit) ? 0 : val + 1;
78    if (val != gpuData[5]) { std::printf("atomicInc failed\n"); return false; }
max (test_web_cs_atomics.cu:69)
59    val = 0;
60    for (int i = 0; i < len; ++i) val -= 10;
61    if (val != gpuData[1]) { std::printf("atomicSub failed\n"); return false; }
62
63    bool found = false;
64    for (int i = 0; i < len; ++i) if (i == gpuData[2]) { found = true; break; }
65    if (!found) { std::printf("atomicExch failed\n"); return false; }
66
67    val = -(1 << 8);
68    for (int i = 0; i < len; ++i) val = max(val, i);
69    if (val != gpuData[3]) { std::printf("atomicMax failed\n"); return false; }
70
71    val = 1 << 8;
72    for (int i = 0; i < len; ++i) val = min(val, i);
73    if (val != gpuData[4]) { std::printf("atomicMin failed\n"); return false; }
74
75    int limit = 17;
76    val = 0;
77    for (int i = 0; i < len; ++i) val = (val >= limit) ? 0 : val + 1;
78    if (val != gpuData[5]) { std::printf("atomicInc failed\n"); return false; }
max (test_web_cs_atomics.cu:69)
59    val = 0;
60    for (int i = 0; i < len; ++i) val -= 10;
61    if (val != gpuData[1]) { std::printf("atomicSub failed\n"); return false; }
62
63    bool found = false;
64    for (int i = 0; i < len; ++i) if (i == gpuData[2]) { found = true; break; }
65    if (!found) { std::printf("atomicExch failed\n"); return false; }
66
67    val = -(1 << 8);
68    for (int i = 0; i < len; ++i) val = max(val, i);
69    if (val != gpuData[3]) { std::printf("atomicMax failed\n"); return false; }
70
71    val = 1 << 8;
72    for (int i = 0; i < len; ++i) val = min(val, i);
73    if (val != gpuData[4]) { std::printf("atomicMin failed\n"); return false; }
74
75    int limit = 17;
76    val = 0;
77    for (int i = 0; i < len; ++i) val = (val >= limit) ? 0 : val + 1;
78    if (val != gpuData[5]) { std::printf("atomicInc failed\n"); return false; }
max (test_web_cs_atomics.cu:69)
59    val = 0;
60    for (int i = 0; i < len; ++i) val -= 10;
61    if (val != gpuData[1]) { std::printf("atomicSub failed\n"); return false; }
62
63    bool found = false;
64    for (int i = 0; i < len; ++i) if (i == gpuData[2]) { found = true; break; }
65    if (!found) { std::printf("atomicExch failed\n"); return false; }
66
67    val = -(1 << 8);
68    for (int i = 0; i < len; ++i) val = max(val, i);
69    if (val != gpuData[3]) { std::printf("atomicMax failed\n"); return false; }
70
71    val = 1 << 8;
72    for (int i = 0; i < len; ++i) val = min(val, i);
73    if (val != gpuData[4]) { std::printf("atomicMin failed\n"); return false; }
74
75    int limit = 17;
76    val = 0;
77    for (int i = 0; i < len; ++i) val = (val >= limit) ? 0 : val + 1;
78    if (val != gpuData[5]) { std::printf("atomicInc failed\n"); return false; }
max (test_web_cs_atomics.cu:69)
59    val = 0;
60    for (int i = 0; i < len; ++i) val -= 10;
61    if (val != gpuData[1]) { std::printf("atomicSub failed\n"); return false; }
62
63    bool found = false;
64    for (int i = 0; i < len; ++i) if (i == gpuData[2]) { found = true; break; }
65    if (!found) { std::printf("atomicExch failed\n"); return false; }
66
67    val = -(1 << 8);
68    for (int i = 0; i < len; ++i) val = max(val, i);
69    if (val != gpuData[3]) { std::printf("atomicMax failed\n"); return false; }
70
71    val = 1 << 8;
72    for (int i = 0; i < len; ++i) val = min(val, i);
73    if (val != gpuData[4]) { std::printf("atomicMin failed\n"); return false; }
74
75    int limit = 17;
76    val = 0;
77    for (int i = 0; i < len; ++i) val = (val >= limit) ? 0 : val + 1;
78    if (val != gpuData[5]) { std::printf("atomicInc failed\n"); return false; }
min (test_web_cs_atomics.cu:73)
63    bool found = false;
64    for (int i = 0; i < len; ++i) if (i == gpuData[2]) { found = true; break; }
65    if (!found) { std::printf("atomicExch failed\n"); return false; }
66
67    val = -(1 << 8);
68    for (int i = 0; i < len; ++i) val = max(val, i);
69    if (val != gpuData[3]) { std::printf("atomicMax failed\n"); return false; }
70
71    val = 1 << 8;
72    for (int i = 0; i < len; ++i) val = min(val, i);
73    if (val != gpuData[4]) { std::printf("atomicMin failed\n"); return false; }
74
75    int limit = 17;
76    val = 0;
77    for (int i = 0; i < len; ++i) val = (val >= limit) ? 0 : val + 1;
78    if (val != gpuData[5]) { std::printf("atomicInc failed\n"); return false; }
79
80    limit = 137;
81    val = 0;
82    for (int i = 0; i < len; ++i) val = ((val == 0) || (val > limit)) ? limit : val - 1;
83    if (val != gpuData[6]) { std::printf("atomicDec failed\n"); return false; }
min (test_web_cs_atomics.cu:73)
63    bool found = false;
64    for (int i = 0; i < len; ++i) if (i == gpuData[2]) { found = true; break; }
65    if (!found) { std::printf("atomicExch failed\n"); return false; }
66
67    val = -(1 << 8);
68    for (int i = 0; i < len; ++i) val = max(val, i);
69    if (val != gpuData[3]) { std::printf("atomicMax failed\n"); return false; }
70
71    val = 1 << 8;
72    for (int i = 0; i < len; ++i) val = min(val, i);
73    if (val != gpuData[4]) { std::printf("atomicMin failed\n"); return false; }
74
75    int limit = 17;
76    val = 0;
77    for (int i = 0; i < len; ++i) val = (val >= limit) ? 0 : val + 1;
78    if (val != gpuData[5]) { std::printf("atomicInc failed\n"); return false; }
79
80    limit = 137;
81    val = 0;
82    for (int i = 0; i < len; ++i) val = ((val == 0) || (val > limit)) ? limit : val - 1;
83    if (val != gpuData[6]) { std::printf("atomicDec failed\n"); return false; }
min (test_web_cs_atomics.cu:73)
63    bool found = false;
64    for (int i = 0; i < len; ++i) if (i == gpuData[2]) { found = true; break; }
65    if (!found) { std::printf("atomicExch failed\n"); return false; }
66
67    val = -(1 << 8);
68    for (int i = 0; i < len; ++i) val = max(val, i);
69    if (val != gpuData[3]) { std::printf("atomicMax failed\n"); return false; }
70
71    val = 1 << 8;
72    for (int i = 0; i < len; ++i) val = min(val, i);
73    if (val != gpuData[4]) { std::printf("atomicMin failed\n"); return false; }
74
75    int limit = 17;
76    val = 0;
77    for (int i = 0; i < len; ++i) val = (val >= limit) ? 0 : val + 1;
78    if (val != gpuData[5]) { std::printf("atomicInc failed\n"); return false; }
79
80    limit = 137;
81    val = 0;
82    for (int i = 0; i < len; ++i) val = ((val == 0) || (val > limit)) ? limit : val - 1;
83    if (val != gpuData[6]) { std::printf("atomicDec failed\n"); return false; }
min (test_web_cs_atomics.cu:73)
63    bool found = false;
64    for (int i = 0; i < len; ++i) if (i == gpuData[2]) { found = true; break; }
65    if (!found) { std::printf("atomicExch failed\n"); return false; }
66
67    val = -(1 << 8);
68    for (int i = 0; i < len; ++i) val = max(val, i);
69    if (val != gpuData[3]) { std::printf("atomicMax failed\n"); return false; }
70
71    val = 1 << 8;
72    for (int i = 0; i < len; ++i) val = min(val, i);
73    if (val != gpuData[4]) { std::printf("atomicMin failed\n"); return false; }
74
75    int limit = 17;
76    val = 0;
77    for (int i = 0; i < len; ++i) val = (val >= limit) ? 0 : val + 1;
78    if (val != gpuData[5]) { std::printf("atomicInc failed\n"); return false; }
79
80    limit = 137;
81    val = 0;
82    for (int i = 0; i < len; ++i) val = ((val == 0) || (val > limit)) ? limit : val - 1;
83    if (val != gpuData[6]) { std::printf("atomicDec failed\n"); return false; }
min (test_web_cs_atomics.cu:73)
63    bool found = false;
64    for (int i = 0; i < len; ++i) if (i == gpuData[2]) { found = true; break; }
65    if (!found) { std::printf("atomicExch failed\n"); return false; }
66
67    val = -(1 << 8);
68    for (int i = 0; i < len; ++i) val = max(val, i);
69    if (val != gpuData[3]) { std::printf("atomicMax failed\n"); return false; }
70
71    val = 1 << 8;
72    for (int i = 0; i < len; ++i) val = min(val, i);
73    if (val != gpuData[4]) { std::printf("atomicMin failed\n"); return false; }
74
75    int limit = 17;
76    val = 0;
77    for (int i = 0; i < len; ++i) val = (val >= limit) ? 0 : val + 1;
78    if (val != gpuData[5]) { std::printf("atomicInc failed\n"); return false; }
79
80    limit = 137;
81    val = 0;
82    for (int i = 0; i < len; ++i) val = ((val == 0) || (val > limit)) ? limit : val - 1;
83    if (val != gpuData[6]) { std::printf("atomicDec failed\n"); return false; }
min (test_web_cs_atomics.cu:73)
63    bool found = false;
64    for (int i = 0; i < len; ++i) if (i == gpuData[2]) { found = true; break; }
65    if (!found) { std::printf("atomicExch failed\n"); return false; }
66
67    val = -(1 << 8);
68    for (int i = 0; i < len; ++i) val = max(val, i);
69    if (val != gpuData[3]) { std::printf("atomicMax failed\n"); return false; }
70
71    val = 1 << 8;
72    for (int i = 0; i < len; ++i) val = min(val, i);
73    if (val != gpuData[4]) { std::printf("atomicMin failed\n"); return false; }
74
75    int limit = 17;
76    val = 0;
77    for (int i = 0; i < len; ++i) val = (val >= limit) ? 0 : val + 1;
78    if (val != gpuData[5]) { std::printf("atomicInc failed\n"); return false; }
79
80    limit = 137;
81    val = 0;
82    for (int i = 0; i < len; ++i) val = ((val == 0) || (val > limit)) ? limit : val - 1;
83    if (val != gpuData[6]) { std::printf("atomicDec failed\n"); return false; }
min (test_web_cs_atomics.cu:73)
63    bool found = false;
64    for (int i = 0; i < len; ++i) if (i == gpuData[2]) { found = true; break; }
65    if (!found) { std::printf("atomicExch failed\n"); return false; }
66
67    val = -(1 << 8);
68    for (int i = 0; i < len; ++i) val = max(val, i);
69    if (val != gpuData[3]) { std::printf("atomicMax failed\n"); return false; }
70
71    val = 1 << 8;
72    for (int i = 0; i < len; ++i) val = min(val, i);
73    if (val != gpuData[4]) { std::printf("atomicMin failed\n"); return false; }
74
75    int limit = 17;
76    val = 0;
77    for (int i = 0; i < len; ++i) val = (val >= limit) ? 0 : val + 1;
78    if (val != gpuData[5]) { std::printf("atomicInc failed\n"); return false; }
79
80    limit = 137;
81    val = 0;
82    for (int i = 0; i < len; ++i) val = ((val == 0) || (val > limit)) ? limit : val - 1;
83    if (val != gpuData[6]) { std::printf("atomicDec failed\n"); return false; }
umin (test_web_cs_mergesort.cu:59)
49                     (msg), static_cast<int>(_e), __FILE__, __LINE__);         \
50        std::exit(EXIT_FAILURE);                                               \
51    } } while (0)
52
53// verbatim device helpers from the source repo
54template <uint sortDir> static inline __device__ uint binarySearchInclusive(uint val, uint *data, uint L, uint stride)
55{
56    if (L == 0) return 0;
57    uint pos = 0;
58    for (; stride > 0; stride >>= 1) {
59        uint newPos = umin(pos + stride, L);
60        if ((sortDir && (data[newPos - 1] <= val)) || (!sortDir && (data[newPos - 1] >= val)))
61            pos = newPos;
62    }
63    return pos;
64}
65
66template <uint sortDir> static inline __device__ uint binarySearchExclusive(uint val, uint *data, uint L, uint stride)
67{
68    if (L == 0) return 0;
69    uint pos = 0;
70    for (; stride > 0; stride >>= 1) {
On this page