小编Eth*_* L.的帖子

ENQCMD和MOVDIR64B是英特尔 DSA 中的两条指令。

MOVDIR64B从源内存地址读取 64 字节，并对目标地址执行 64 字节直接存储操作。该ENQCMD指令允许软件将命令写入队列寄存器，这些寄存器是使用内存映射 I/O (MMIO) 访问的特殊设备寄存器。

我的问题是——设计这两条指令的目的是什么？

根据我的理解，设置内存映射IO区域（寄存器）需要操作系统的支持，即设备驱动程序。设置MMIO区域后，我们可以使用write()系统调用来访问它，这也是在设备驱动程序中实现的。对于一般架构，Linux 支持iowrite64()一次写入 8 字节值。因此，如果我们要写入64字节，需要调用iowrite64()8次。

在英特尔 DSA 的帮助下MOVDIR64B，创建了一个新的 API__iowrite512() -以原子方式写入 64 字节。

我同意后一种至少比前一种更有效，但我对传输数据所需的时间感到困惑。

考虑以下情况：如果给定一个支持MOVDIR64B和 的设备（Intel DSA） ENQCMD，假设我们想要将 64 字节数据从内存传输到 MMIO 寄存器。有两种选择：iowrite64()8次（使用循环）；或__iowrite512()一次。后一个会比前一个快8倍吗？

我的想法是8倍差别的可能性较小，但后者会更快。我可以知道它会多快吗？它在任何地方都有记录吗？我没有 Intel DSA，所以我不知道如何测试它。

除此之外，还有什么好处呢ENQCMD？会不会被分解成几个微操作？如果是的话，那么微操作有哪些作用呢ENQCMD？

假设我们有一个调用一些函数的内核，例如：

__device__ int fib(int n) {
    if (n == 0 || n == 1) {
        return n;
    } else {
        int x = fib(n-1);
        int y = fib(n-2);
        return x + y;
    }
    return -1;
}

__global__ void fib_kernel(int* n, int *ret) {
    *ret = fib(*n);
}

内核fib_kernel将调用 function fib()，该函数在内部将调用两个fib()函数。假设GPU有80个SM，我们正好启动80个线程来进行计算，并传入n10个。我知道会有大量的重复计算，这违反了数据并行性的思想，但我想更好地理解线程的堆栈管理。

根据Cuda PTX 的文档，它声明如下：

GPU 维护每个线程的执行状态，包括程序计数器和调用堆栈

1. 堆栈位于本地内存中。当线程执行内核时，它们的行为是否与CPU中的调用约定一样？换句话说，是不是对于每个线程，对应的栈都会动态增长和收缩呢？

2. 每个线程的堆栈都是私有的，其他线程无法访问。有没有一种方法可以手动检测编译器/驱动程序，以便堆栈分配在全局内存中，而不是本地内存中？

3. 有没有一种方法可以让线程获取当前的程序计数器、帧指针值？我认为它们存储在一些特定的寄存器中，但 PTX 文档没有提供访问这些寄存器的方法。我可以知道我必须修改什么（例如驱动程序或编译器）才能获取这些寄存器吗？

4. 如果我们将输入增加到fib(n)10000，很可能会导致堆栈溢出，有办法处理吗？问题2的答案或许可以解决这个问题。任何其他想法将不胜感激。