logo标签列表

cudaSetDevice()函数对CUDA设备的上下文堆栈有什么影响？

cudacuda-contextcuda-driver
5
假设我有一个与设备 i 相关联的活动CUDA上下文，现在我调用 cudaSetDevice（i）。 会发生什么？
  1. 什么也不会发生？
  2. 主要上下文替换堆栈顶部？
  3. 将主要上下文推送到堆栈上？
实际上这似乎是不一致的。 我已经编写了这个程序，在单个设备的计算机上运行：
#include <cuda.h>
#include <cuda_runtime_api.h>
#include <cassert>
#include <iostream>

int main()
{
        CUcontext ctx1, primary;
        cuInit(0);
        auto status = cuCtxCreate(&ctx1, 0, 0);
        assert (status == (CUresult) cudaSuccess);
        cuCtxPushCurrent(ctx1);
        status = cudaSetDevice(0);
        assert (status == cudaSuccess);
        void* ptr1;
        void* ptr2;
        cudaMalloc(&ptr1, 1024);
        assert (status == cudaSuccess);
        cuCtxGetCurrent(&primary);
        assert (status == (CUresult) cudaSuccess);
        assert(primary != ctx1);
        status = cuCtxPushCurrent(ctx1);
        assert (status == (CUresult) cudaSuccess);
        cudaMalloc(&ptr2, 1024);
        assert (status == (CUresult) cudaSuccess);
        cudaSetDevice(0);
        assert (status == (CUresult) cudaSuccess);
        int i = 0;
        while (true) {
                status = cuCtxPopCurrent(&primary);
                if (status != (CUresult) cudaSuccess) { break; }
                std::cout << "Next context on stack (" << i++ << ") is " << (void*) primary << '\n';
        }
}

我得到了以下输出：

context ctx1 is 0x563ec6225e30
primary context is 0x563ec61f5490
Next context on stack (0) is 0x563ec61f5490
Next context on stack (1) is 0x563ec61f5490
Next context on stack(2) is 0x563ec6225e3

这似乎是行为 有时 是替换，而有时是推送。

发生了什么？

1个回答
6

简述：根据您提供的代码，在您特定的两种用法中，似乎cudaSetDevice（）正在替换堆栈顶部的上下文。

让我们稍微修改您的代码，然后看看我们可以推断出代码中每个API调用对上下文堆栈的影响：

$ cat t1759.cu
#include <cuda.h>
#include <cuda_runtime_api.h>
#include <cassert>
#include <iostream>
void check(int j, CUcontext ctx1, CUcontext ctx2){
  CUcontext ctx0;
  int i = 0;
  while (true) {
                auto status = cuCtxPopCurrent(&ctx0);
                if (status != CUDA_SUCCESS) { break; }
                if (ctx0 == ctx1) std::cout << j << ":Next context on stack (" << i++ << ") is ctx1:" << (void*) ctx0 << '\n';
                else if (ctx0 == ctx2) std::cout << j << ":Next context on stack (" << i++ << ") is ctx2:" << (void*) ctx0 << '\n';
                else std::cout << j << ":Next context on stack (" << i++ << ") is unknown:" << (void*) ctx0 << '\n';
  }
}
void runtest(int i)
{
        CUcontext ctx1, primary = NULL;
        cuInit(0);
        auto dstatus = cuCtxCreate(&ctx1, 0, 0);    // checkpoint 1
        assert (dstatus == CUDA_SUCCESS);
        if (i == 1) {check(i,ctx1,primary); return;}// checkpoint 1
        dstatus = cuCtxPushCurrent(ctx1);           // checkpoint 2
        assert (dstatus == CUDA_SUCCESS);
        if (i == 2) {check(i,ctx1,primary); return;}// checkpoint 2
        auto rstatus = cudaSetDevice(0);            // checkpoint 3
        assert (rstatus == cudaSuccess);
        if (i == 3) {check(i,ctx1,primary); return;}// checkpoint 3
        void* ptr1;
        void* ptr2;
        rstatus = cudaMalloc(&ptr1, 1024);          // checkpoint 4
        assert (rstatus == cudaSuccess);
        if (i == 4) {check(i,ctx1,primary); return;}// checkpoint 4
        dstatus = cuCtxGetCurrent(&primary);        // checkpoint 5
        assert (dstatus == CUDA_SUCCESS);
        assert(primary != ctx1);
        if (i == 5) {check(i,ctx1,primary); return;}// checkpoint 5
        dstatus = cuCtxPushCurrent(ctx1);           // checkpoint 6
        assert (dstatus == CUDA_SUCCESS);
        if (i == 6) {check(i,ctx1,primary); return;}// checkpoint 6
        rstatus = cudaMalloc(&ptr2, 1024);          // checkpoint 7
        assert (rstatus == cudaSuccess);
        if (i == 7) {check(i,ctx1,primary); return;}// checkpoint 7
        rstatus = cudaSetDevice(0);                 // checkpoint 8
        assert (rstatus == cudaSuccess);
        if (i == 8) {check(i,ctx1,primary); return;}// checkpoint 8
        return;
}

int main(){
        for (int i = 1; i < 9; i++){
          cudaDeviceReset();
          runtest(i);}
}
$ nvcc -o t1759 t1759.cu -lcuda -std=c++11
$ ./t1759
1:Next context on stack (0) is ctx1:0x11087e0
2:Next context on stack (0) is ctx1:0x1741160
2:Next context on stack (1) is ctx1:0x1741160
3:Next context on stack (0) is unknown:0x10dc520
3:Next context on stack (1) is ctx1:0x1c5aa70
4:Next context on stack (0) is unknown:0x10dc520
4:Next context on stack (1) is ctx1:0x23eaa00
5:Next context on stack (0) is ctx2:0x10dc520
5:Next context on stack (1) is ctx1:0x32caf30
6:Next context on stack (0) is ctx1:0x3a44ed0
6:Next context on stack (1) is ctx2:0x10dc520
6:Next context on stack (2) is ctx1:0x3a44ed0
7:Next context on stack (0) is ctx1:0x41cfd90
7:Next context on stack (1) is ctx2:0x10dc520
7:Next context on stack (2) is ctx1:0x41cfd90
8:Next context on stack (0) is ctx2:0x10dc520
8:Next context on stack (1) is ctx2:0x10dc520
8:Next context on stack (2) is ctx1:0x4959c70
$

基于上述内容，当我们在代码中进行每个API调用时：
1.
        auto dstatus = cuCtxCreate(&ctx1, 0, 0);    // checkpoint 1
1:Next context on stack (0) is ctx1:0x11087e0

正如此处所述，上下文创建还会将新创建的上下文推送到堆栈上。

2.

        dstatus = cuCtxPushCurrent(ctx1);           // checkpoint 2
2:Next context on stack (0) is ctx1:0x1741160
2:Next context on stack (1) is ctx1:0x1741160

没有什么奇怪的，将相同的上下文推入栈中会为其创建另一个栈条目。
3.
        auto rstatus = cudaSetDevice(0);            // checkpoint 3
3:Next context on stack (0) is unknown:0x10dc520
3:Next context on stack (1) is ctx1:0x1c5aa70
cudaSetDevice()函数调用后，将栈顶替换为一个“未知”的上下文（此时尚未检索到“其他”上下文的句柄值）。
        rstatus = cudaMalloc(&ptr1, 1024);          // checkpoint 4
4:Next context on stack (0) is unknown:0x10dc520
4:Next context on stack (1) is ctx1:0x23eaa00

由于此调用，堆栈配置无变化。

5.

        dstatus = cuCtxGetCurrent(&primary);        // checkpoint 5
5:Next context on stack (0) is ctx2:0x10dc520
5:Next context on stack (1) is ctx1:0x32caf30

由于这个调用，并没有改变堆栈配置，但我们现在知道栈顶上下文是当前上下文（并且我们可以推测它是主要上下文）。
6.
        dstatus = cuCtxPushCurrent(ctx1);           // checkpoint 6
6:Next context on stack (0) is ctx1:0x3a44ed0
6:Next context on stack (1) is ctx2:0x10dc520
6:Next context on stack (2) is ctx1:0x3a44ed0

没什么意外。我们将ctx1推入栈中，因此栈有3个条目，第一个是驱动程序API创建的上下文，接下来两个条目与步骤5中的堆栈配置相同，只是向下移动了一个堆栈位置。
        rstatus = cudaMalloc(&ptr2, 1024);          // checkpoint 7
7:Next context on stack (0) is ctx1:0x41cfd90
7:Next context on stack (1) is ctx2:0x10dc520
7:Next context on stack (2) is ctx1:0x41cfd90

再次强调，此调用对堆栈配置没有影响。

8.

        rstatus = cudaSetDevice(0);                 // checkpoint 8
8:Next context on stack (0) is ctx2:0x10dc520
8:Next context on stack (1) is ctx2:0x10dc520
8:Next context on stack (2) is ctx1:0x4959c70

我们再次看到这里的行为是，cudaSetDevice()调用已经用主上下文替换了堆栈顶部的上下文。

从您的测试代码中得出的结论是，在您的代码中与各种运行时和驱动程序API调用混合使用时，cudaSetDevice()调用的行为没有任何不一致性。

在我看来，这种编程范例是疯狂的。我无法想象为什么您要以这种方式混合使用驱动程序API和运行时API代码。

哦，不，我不想这样做，那只是一个人工的例子，我只是想了解行为，这样当我将我的API包装器扩展到驱动程序API时，我就不会因错误的假设而搞砸事情。 - einpoklum
现在，我有一段代码可以“推送”和“弹出”当前设备 - 假设只使用运行时API。现在您已经澄清了cudaSetDevice()的行为，我可以更改该代码以执行以下操作：1.查看当前上下文。2.将其保存在侧面。3.使用运行时API将cudaSetDevice()设置为要使用的设备4.完成我的工作。5.将cuCtxSetCurrent()设置为我保存的上下文。 - einpoklum
通过编辑我的回答，您做出了一个我不舒服的声明，并有效地将该声明归因于我。除非我将它们标记为社区wiki，否则最好不要编辑我的回答。我对此感到不舒服，我建议您不要在未来这样擅自修改我的回答。我喜欢就我所观察到的事情发表声明。就让它保持这样吧。如果您需要进一步的信息，请通过developer.nvidia.com上的错误报告门户提出请求。如果您需要澄清，请使用评论。 - Robert Crovella
此外，我对你在评论中列出的5个步骤不做任何声明。 - Robert Crovella
我会尊重您的建议和请求。 - einpoklum
我已经在我的结合了运行时和驱动程序API包装器的现代C++实现中使用了这个见解。谢谢（如果您对此有反馈，将不胜感激）。 - einpoklum
网页内容由stack overflow 提供, 点击上面的
可以查看英文原文，
原文链接