CUDA中的一些数学运算答案

【问题标题】：some mathematical operations in CUDACUDA中的一些数学运算
【发布时间】：2013-07-12 17:42:25
【问题描述】：

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

【问题讨论】：

标签： c cuda gpgpu

【解决方案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 运算符 (&&) 的使用是否符合您对“数学运算”的定义。

【讨论】：

是的，我知道，这是一个普遍的问题。我只是想在 cuda 应用程序中使用它时进行标记。我刚刚测试了表达式，当我输入 2 时它给出 0。我想要 1 代替 0。
我更正了你的表达。右移后用 A 代替 1
我无法编辑您的帖子。请编辑它，我会接受它作为答案。太感谢了。你刚刚救了我。
我的表达方式是这样的。请测试一下。右移后不需要用 A 代替 1。请注意，我更新了表达式以添加额外的括号。请重新测试。
您可以尝试的其他一些方法是x = min (x, 1)、x = (x > 1) ? 1 : x 和x = x - (x >> 1)。

【解决方案2】：

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

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

但是，如果您想避免分支以最大限度地提高速度：自 PTX ISA 1.0 以来就有一条 min 指令。（参见 PTX ISA 3.1 手册中的 Tab. 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 的 Sec. 6.5.3.3 Par. 5 和 Sec. 6.5.9 Par. 3. Afair这个保证也适用于 CUDA。）

【讨论】：