c - CUDA中的一些数学运算

Question

我有一个包含 0,1 和 2 的二维矩阵。我正在编写一个 cuda 内核，其中线程数等于矩阵大小，每个线程将对矩阵的每个元素进行操作。现在，我需要可以保持 0 和 1 不变的数学运算，但会将 2 转换为 1。这是一个数学运算，没有任何 if-else，它将执行以下转换：0 -> 0; 1 -> 1; 2 -> 1。有没有可能使用数学运算符进行上述转换的方法。任何帮助将不胜感激。谢谢你。

Answer 1

这不是一个 cuda 问题。

int A;
// set A to 0, 1, or 2
int a = (A + (A>>1)) & 1;
// a is now 0 if A is 0, or 1 if A is 1 or 2

或作为宏：

#define fix01(x) ((x+(x>>1))&1)

int a = fix01(A);

这似乎也有效：

#define fix01(x) ((x&&1)&1)

我不知道布尔 AND 运算符 ( &&) 的使用是否符合您对“数学运算”的定义。

Answer 2

由于问题是关于“数学”函数，我建议使用以下二阶多项式：

int f(int x) { return ((3-x)*x)/2; }

但是，如果您想避免分支以最大限度地提高速度：自 PTX ISA 1.0 以来就有一条 min 指令。（请参阅 PTX ISA 3.1 手册中的表 36。）因此以下 CUDA 代码

__global__ void test(int *x, int *y)
{
    *y = *x <= 1 ? *x : 1;
}

在我的测试中编译为以下 PTX 汇编器（只是从 CUDA 5 调用 nvcc，没有任何拱选项）

    code for sm_10
            Function : _Z4testPiS_
    /*0000*/     /*0x1000c8010423c780*/     MOV R0, g [0x4];
    /*0008*/     /*0xd00e000580c00780*/     GLD.U32 R1, global14 [R0];
    /*0010*/     /*0x1000cc010423c780*/     MOV R0, g [0x6];
    /*0018*/     /*0x30800205ac400780*/     IMIN.S32 R1, R1, c [0x1] [0x0];
    /*0020*/     /*0xd00e0005a0c00781*/     GST.U32 global14 [R0], R1;

因此，使用条件 ?: 的 min() 实现实际上编译为单个 IMIN.S32 PTX 指令，没有任何分支。因此，我建议将其用于任何实际应用程序：

int f(int x) { return x <= 1 ? x : 1; }

但回到只使用非分支操作的问题：

在 C 中获得此结果的另一种形式是使用两个非运算符：

int f(int x) { return !!x; }

或者简单地与零比较：

int f(int x) { return x != 0; }

（! 和 != 的结果保证为 0 或 1，比较 C99 标准 ISO/IEC 9899:1999 的第 6.5.3.3 节第 5 节和第 6.5.9 节第 3 节。公平这一保证也在 CUDA 中成立。）