optimization - 有效地将无符号值除以 2 的幂，四舍五入 - 在 CUDA 中

Question

我只是在读:

Efficiently dividing unsigned value by a power of two, rounding up

我想知道在 CUDA 中最快的方法是什么。当然，“快速”是指吞吐量方面(该问题还解决了后续调用相互依赖的情况)。

对于 lg()该问题中提到的函数(除数的基数为 2 的对数)，假设我们有:

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 中没有 find-first-sert，到目前为止所有 nVIDIA GPU 的指令集中也没有，让我们替换一下 lg()具有以下内容:

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); }

我们现在需要实现

template <typename T> T div_by_power_of_2_rounding_up(T p, T q);

... 为 T = uint32_t和 T = uint64_t . ( p 是被除数，q 是除数)。

备注:

与原始问题一样，我们可能不会假设 p <= std::numeric_limits<T>::max() - q或者那个 p > 0 - 这会使各种有趣的替代方案崩溃:-)

0 不是 2 的幂，所以我们可以假设 q != 0 .

我意识到 32 位和 64 位的解决方案可能不同；我对前者更感兴趣，但也对后者感兴趣。

让我们专注于 Maxwell 和 Pascal 芯片。

Answer 1

使用漏斗移位，可能的 32 位策略是进行 33 位移位(基本上)保留加法的进位，以便在移位之前完成，例如:(未测试)

unsigned sum = dividend + mask;
unsigned result = __funnelshift_r(sum, sum < mask, log_2_of_divisor);

由@einpoklum 编辑 :

使用@RobertCrovella 的程序进行测试，似乎工作正常。 SM_61 的测试内核 PTX 是:

    .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 */