从常量存储器统一加载 128 位数据

Question

给定一个 CUDA 向量类型 int4，我如何从常量内存中加载 128 位数据。

这似乎不起作用：

#include <stdio.h>
#include <cuda.h>

__constant__ int constant_mem[4];
__global__ void kernel(){
    int4 vec;
    vec = constant_mem[0];
}
int main(void){return 0;}

在第七行，我尝试将常量内存中的所有 4 个整数值加载到 128 位向量类型中。此操作导致以下编译错误：

vectest.cu(7): error: no operator "=" matches these operands
            operand types are: int4 = int

另外，是否可以直接访问向量类型而无需强制转换，如下所示：

int data = vec[0];

PTX 程序集中的 Switch 语句：

    @%p1 bra    BB1_55;

    setp.eq.s32     %p26, %r1, 1;
    @%p26 bra   BB1_54;

    setp.eq.s32     %p27, %r1, 2;
    @%p27 bra   BB1_53;

    setp.ne.s32     %p28, %r1, 3;
    @%p28 bra   BB1_55;

    mov.u32     %r961, %r61;
    bra.uni     BB1_56;

BB1_53:
    mov.u32     %r961, %r60;
    bra.uni     BB1_56;

BB1_54:
    mov.u32     %r961, %r59;
    bra.uni     BB1_56;

BB1_55:
    mov.u32     %r961, %r58;

BB1_56:

Answer 1

在第一种情况下，强制转换可能是最简单的解决方案，所以像这样：

__constant__ int constant_mem[4];
__global__ void kernel(){
    int4 vec = * reinterpret_cast<int4 *>(&constant_mem);
}

（在浏览器中编写的免责声明，未经编译或测试，使用风险自负）

使用 C++ reinterpret_cast 运算符将强制编译器发出 128 位加载指令。

在第二种情况下，听起来您想使用 128 位内存事务直接寻址存储在 128 位向量类型数组中的 32 位字。这需要一些辅助函数，可能类似于：

__inline__ __device__ int fetch4(const int4 val, const int n)
{
     (void) val.x; (void) val.y; (void) val.z; (void) val.w;
     switch(n) {
         case 3:
            return val.w;
         case 2: 
            return val.z;
         case 1:
            return val.y;
         case 0:
         default:
            return val.x;
    }
}

__device__ int index4(const int4 * array, const int n)
{
    int div = n / 4;
    int mod = n - (div * 4);

    int4 val = array[div]; // 128 bit load here

    return fetch4(val, mod);
}

__constant__ int constant_mem[128];
__global__ void kernel(){
    int val = index4(constant_mem, threadIdx.x);
}

（在浏览器中编写的免责声明，未经编译或测试，使用风险自负）

在这里，我们通过读取整个 int4 值并解析其内容来强制执行 128 位事务（强制转换为 void 是旧版本的 open64 编译器所必需的咒语，如果它认为成员未使用，则倾向于优化向量加载）。执行索引需要一些 IOP 开销，但如果生成的事务的负载带宽更高，它们可能是值得的。 switch 语句可能是使用条件执行编译的，所以不应该有分支分歧惩罚。请注意，非常随机地访问 int4 值数组可能会浪费大量带宽并导致扭曲序列化。这样做可能会对性能产生很大的负面影响。