一种可能的简单方法是:
$ cat t105.cu
#include <cstdio>
__device__ unsigned r = 0;
template <typename T>
__device__ T pdqru(T p, T q){
T p1 = p + (q-1);
if (sizeof(T) == 8)
q = __ffsll(q);
else
q = __ffs(q);
return (p1<p)?((p>>(q-1))+1) :(p1 >> (q-1));
}
__global__ void test(unsigned p, unsigned q){
#ifdef USE_DISPLAY
unsigned q2 = 16;
unsigned z = 0;
unsigned l = 1U<<31;
printf("result %u/%u = %u\n", p, q, pdqru(p, q));
printf("result %u/%u = %u\n", p, q2, pdqru(p, q2));
printf("result %u/%u = %u\n", p, z, pdqru(p, z));
printf("result %u/%u = %u\n", z, q, pdqru(z, q));
printf("result %u/%u = %u\n", l, q, pdqru(l, q));
printf("result %u/%u = %u\n", q, l, pdqru(q, l));
printf("result %u/%u = %u\n", l, l, pdqru(l, l));
printf("result %u/%u = %u\n", q, q, pdqru(q, q));
#else
r = pdqru(p, q);
#endif
}
int main(){
unsigned h_r;
test<<<1,1>>>(32767, 32);
cudaMemcpyFromSymbol(&h_r, r, sizeof(unsigned));
printf("result = %u\n", h_r);
}
$ nvcc -arch=sm_61 -o t105 t105.cu
$ cuobjdump -sass ./t105
Fatbin elf code:
================
arch = sm_61
code version = [1,7]
producer = <unknown>
host = linux
compile_size = 64bit
code for sm_61
Fatbin elf code:
================
arch = sm_61
code version = [1,7]
producer = cuda
host = linux
compile_size = 64bit
code for sm_61
Function : _Z4testjj
.headerflags @"EF_CUDA_SM61 EF_CUDA_PTX_SM(EF_CUDA_SM61)"
MOV R1, c[0x0][0x20];
IADD R0, RZ, -c[0x0][0x144];
LOP.AND R0, R0, c[0x0][0x144];
FLO.U32 R2, R0;
MOV R0, c[0x0][0x144];
IADD32I R3, -R2, 0x1f;
IADD32I R0, R0, -0x1;
MOV R2, c[0x0][0x140];
IADD32I R4, -R3, 0x1f;
IADD R5, R0, c[0x0][0x140];
ISETP.LT.U32.AND P0, PT, R5, R0, PT;
SHR.U32 R0, R2, R4;
IADD32I R0, R0, 0x1;
MOV32I R2, 0x0;
MOV32I R3, 0x0;
@!P0 SHR.U32 R0, R5, R4;
STG.E [R2], R0;
EXIT;
BRA 0xc0;
NOP;
NOP;
NOP;
NOP;
NOP;
..........................
Fatbin ptx code:
================
arch = sm_61
code version = [5,0]
producer = cuda
host = linux
compile_size = 64bit
compressed
$ nvcc -arch=sm_61 -o t105 t105.cu -DUSE_DISPLAY
$ cuda-memcheck ./t105
========= CUDA-MEMCHECK
result 32767/32 = 1024
result 32767/16 = 2048
result 32767/0 = 2048
result 0/32 = 0
result 2147483648/32 = 67108864
result 32/2147483648 = 1
result 2147483648/2147483648 = 1
result 32/32 = 1
result = 0
========= ERROR SUMMARY: 0 errors
$
大约有14个SASS指令适用于32位情况,将答案存储在R0中。对于除以零的情况,它会产生虚假结果。
这个答案的等效汇编代码如下:
$ cat t106.cu
#include <cstdio>
#include <cstdint>
__device__ unsigned r = 0;
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__ T lg(T x) { return find_first_set(x) - 1; }
template <typename T>
__device__ T pdqru(T dividend, T divisor)
{
auto log_2_of_divisor = lg(divisor);
auto mask = divisor - 1;
auto correction_for_rounding_up = ((dividend & mask) + mask) >> log_2_of_divisor;
return (dividend >> log_2_of_divisor) + correction_for_rounding_up;
}
__global__ void test(unsigned p, unsigned q){
#ifdef USE_DISPLAY
unsigned q2 = 16;
unsigned z = 0;
unsigned l = 1U<<31;
printf("result %u/%u = %u\n", p, q, pdqru(p, q));
printf("result %u/%u = %u\n", p, q2, pdqru(p, q2));
printf("result %u/%u = %u\n", p, z, pdqru(p, z));
printf("result %u/%u = %u\n", z, q, pdqru(z, q));
printf("result %u/%u = %u\n", l, q, pdqru(l, q));
printf("result %u/%u = %u\n", q, l, pdqru(q, l));
printf("result %u/%u = %u\n", l, l, pdqru(l, l));
printf("result %u/%u = %u\n", q, q, pdqru(q, q));
#else
r = pdqru(p, q);
#endif
}
int main(){
unsigned h_r;
test<<<1,1>>>(32767, 32);
cudaMemcpyFromSymbol(&h_r, r, sizeof(unsigned));
printf("result = %u\n", h_r);
}
$ nvcc -std=c++11 -arch=sm_61 -o t106 t106.cu
$ cuobjdump -sass t106
Fatbin elf code:
================
arch = sm_61
code version = [1,7]
producer = <unknown>
host = linux
compile_size = 64bit
code for sm_61
Fatbin elf code:
================
arch = sm_61
code version = [1,7]
producer = cuda
host = linux
compile_size = 64bit
code for sm_61
Function : _Z4testjj
.headerflags @"EF_CUDA_SM61 EF_CUDA_PTX_SM(EF_CUDA_SM61)"
MOV R1, c[0x0][0x20];
IADD R0, RZ, -c[0x0][0x144];
MOV R2, c[0x0][0x144];
LOP.AND R0, R0, c[0x0][0x144];
FLO.U32 R3, R0;
IADD32I R0, R2, -0x1;
IADD32I R3, -R3, 0x1f;
LOP.AND R2, R0, c[0x0][0x140];
IADD R2, R0, R2;
IADD32I R0, -R3, 0x1f;
MOV R3, c[0x0][0x140];
MOV32I R6, 0x0;
SHR.U32 R4, R2, R0.reuse;
SHR.U32 R5, R3, R0;
MOV R2, R6;
MOV32I R3, 0x0;
IADD R0, R4, R5;
STG.E [R2], R0;
EXIT;
BRA 0xd0;
NOP;
NOP;
NOP;
NOP;
..........................
Fatbin ptx code:
================
arch = sm_61
code version = [5,0]
producer = cuda
host = linux
compile_size = 64bit
compressed
$
根据我的计算,这似乎比原来多了1条指令。
(被除数+掩码)>>除数的log2
。 - harold