有效地将无符号值除以2的幂，取整-在CUDA中 [英] Efficiently dividing unsigned value by a power of two, rounding up - in CUDA
问题描述
我当时正在阅读：
Efficiently dividing unsigned value by a power of two, rounding up
，我想知道这是什么CUDA中最快的方法。当然，快速是指吞吐量（该问题还解决了彼此之间后续调用的情况）。
and I was wondering what was the fastest way to do this in CUDA. Of course by "fast" I mean in terms of throughput (that question also addressed the case of subsequent calls depending on each other).
对于
For the lg() function mentioned in that question (base-2 logarithm of divisor), suppose we have:
template <typename T> __device__ int find_first_set(T x);
template <> __device__ int find_first_set<uint32_t>(uint32_t x) { return __ffs(x); }
template <> __device__ int find_first_set<uint64_t>(uint64_t x) { return __ffsll(x); }
template <typename T> __device__ int lg(T x) { return find_first_set(x) - 1; }
编辑：既然我已经意识到PTX以及到目前为止所有nVIDIA GPU的指令集中都没有find-first-sert，所以让我们替换 lg（），其中包含以下内容：
Since I've been made aware that there's no find-first-sert in PTX, nor in the instruction set of all nVIDIA GPUs up to this time, let's replace that lg() with the following:
template <typename T> __df__ int population_count(T x);
template <> int population_count<uint32_t>(uint32_t x) { return __popc(x); }
template <> int population_count<uint64_t>(uint64_t x) { return __popcll(x); }
template <typename T>
__device__ int lg_for_power_of_2(T x) { return population_count(x - 1); }
，我们现在需要实施
and we now need to implement
template <typename T> T div_by_power_of_2_rounding_up(T p, T q);
...对于 T = uint32_t 和 T = uint64_t 。 （ p 是股息， q 是除数）。
... for T = uint32_t and T = uint64_t. (p is the dividend, q is the divisor).
注释：
- 与原始问题一样，我们可能不假设
p< = std :: numeric_limits< T> :: max（）-q或p> 0-会折叠各种有趣的替代方法：-） - 0不是2的幂，因此我们可以假设
q！ = 0。 - 我意识到32位和64位解决方案可能有所不同；我对前者更感兴趣，对后者也更感兴趣。
- 让我们关注Maxwell和Pascal芯片。
- As in the original question, we may not assume that
p <= std::numeric_limits<T>::max() - qor thatp > 0- that would collapse the various interesting alternatives :-) - 0 is not a power of 2, so we may assume
q != 0. - I realize solutions might differ for 32-bit and 64-bit; I'm more interested in the former but also in the latter.
- Let's focus on Maxwell and Pascal chips.
推荐答案
可用漏斗移位，一种可能的32位策略是进行33位移位（基本上），以保留加法的进位，以便在移位之前完成，例如
With funnel shifting available, a possible 32 bit strategy is doing a 33bit shift (essentially) preserving the carry of the addition so it be done before the shift, such as this: (not tested)
unsigned sum = dividend + mask;
unsigned result = __funnelshift_r(sum, sum < mask, log_2_of_divisor);
由@einpoklum编辑：
使用@RobertCrovella的程序进行了测试，似乎工作正常。用于SM_61的测试内核PTX为：
Tested using @RobertCrovella's program, seems to work fine. The test kernel PTX for SM_61 is:
.reg .pred %p<2>;
.reg .b32 %r<12>;
ld.param.u32 %r5, [_Z4testjj_param_0];
ld.param.u32 %r6, [_Z4testjj_param_1];
neg.s32 %r7, %r6;
and.b32 %r8, %r6, %r7;
clz.b32 %r9, %r8;
mov.u32 %r10, 31;
sub.s32 %r4, %r10, %r9;
add.s32 %r11, %r6, -1;
add.s32 %r2, %r11, %r5;
setp.lt.u32 %p1, %r2, %r11;
selp.u32 %r3, 1, 0, %p1;
// inline asm
shf.r.wrap.b32 %r1, %r2, %r3, %r4;
// inline asm
st.global.u32 [r], %r1;
ret;
，SASS为：
/*0008*/ MOV R1, c[0x0][0x20]; /* 0x4c98078000870001 */
/*0010*/ MOV R0, c[0x0][0x144]; /* 0x4c98078005170000 */
/*0018*/ IADD R2, RZ, -c[0x0][0x144]; /* 0x4c1100000517ff02 */
/* 0x001c4c00fe4007f1 */
/*0028*/ IADD32I R0, R0, -0x1; /* 0x1c0ffffffff70000 */
/*0030*/ LOP.AND R2, R2, c[0x0][0x144]; /* 0x4c47000005170202 */
/*0038*/ FLO.U32 R2, R2; /* 0x5c30000000270002 */
/* 0x003fd800fe2007e6 */
/*0048*/ IADD R5, R0, c[0x0][0x140]; /* 0x4c10000005070005 */
/*0050*/ ISETP.LT.U32.AND P0, PT, R5, R0, PT; /* 0x5b62038000070507 */
/*0058*/ IADD32I R0, -R2, 0x1f; /* 0x1d00000001f70200 */
/* 0x001fc400fe2007f6 */
/*0068*/ IADD32I R0, -R0, 0x1f; /* 0x1d00000001f70000 */
/*0070*/ SEL R6, RZ, 0x1, !P0; /* 0x38a004000017ff06 */
/*0078*/ MOV32I R2, 0x0; /* 0x010000000007f002 */
/* 0x0003c400fe4007e4 */
/*0088*/ MOV32I R3, 0x0; /* 0x010000000007f003 */
/*0090*/ SHF.R.W R0, R5, R0, R6; /* 0x5cfc030000070500 */
/*0098*/ STG.E [R2], R0; /* 0xeedc200000070200 */
/* 0x001f8000ffe007ff */
/*00a8*/ EXIT; /* 0xe30000000007000f */
/*00b0*/ BRA 0xb0; /* 0xe2400fffff87000f */
/*00b8*/ NOP; /* 0x50b0000000070f00 */
这篇关于有效地将无符号值除以2的幂，取整-在CUDA中的文章就介绍到这了，希望我们推荐的答案对大家有所帮助，也希望大家多多支持IT屋！
