c++ - Thrust::raw_pointer_cast 和多个 GPU，奇怪的行为

Question

我在代码中经常使用推力，因为它是一个很棒的包装器并且提供了非常有用的实用程序，自从添加了异步行为的支持以来我更加确信了。

我的代码使用 cuda Thrust 运行良好，直到我最近尝试在我的应用程序中添加多 GPU 支持。 我经历过烦人的事情

CUDA Runtime API error 77 : an illegal memory access was encountered

我的部分代码以前从未显示过任何边界问题。

我在代码中添加了冗长的内容，似乎我的 Thrust::device_vector 指针地址在执行过程中不断变化，没有明显的原因，在手写内核中生成错误 77。

我可能误解了 UVA 概念及其最终的“副作用”，但我仍然对理解导致指针更新的过程感兴趣。

我无法准确重现我的问题，其中我不使用临时主机变量来存储 cuda 内存指针，而只是在内核包装器调用中需要时动态推力::raw_pointer_cast。

但是我编写了一个小程序，显示了我可能遇到的错误类型，请注意，这并不健壮，您的系统上至少需要有 2 个 GPU 才能运行它:

/********************************************************************************************
** Compile using nvcc ./test.cu -gencode arch=compute_35,code=sm_35 -std=c++11 -o test.exe **
********************************************************************************************/

//Standard Library
#include <iostream>
#include <vector>

//Cuda
#include "cuda_runtime.h"

//Thrust
#include <thrust/device_vector.h>
#include <thrust/functional.h>
#include <thrust/transform.h>

inline void __checkCudaErrors( cudaError err, const char *file, const int line )
{
    if( err != cudaSuccess )
    {
        printf("%s(%i) : CUDA Runtime API error %i : %s \n",file ,line, (int)err, cudaGetErrorString(err) );
    }
};

#define checkCudaErrors(err)    __checkCudaErrors (err, __FILE__, __LINE__)

__global__ void write_memory( float* buf, float value )
{
    printf("GPU TALK: Raw pointer is %p \n",buf);
    buf[0] = value;
}

int main( int argc, char* argv[] )
{
    //declare a vector of vector
    std::vector<thrust::device_vector<float> > v;
    float test;
    float* tmp;

    //Initialize first vector on GPU 0
    cudaSetDevice( 0 );
    v.emplace_back( 65536, 1.0f );
    tmp = thrust::raw_pointer_cast( v.at(0).data() );
    std::cout << " Host TALK: Raw pointer of vector 0 at step 0 " << (void*)tmp << std::endl;

    //Try to use it raw pointer
    write_memory<<<1,1,0,0>>>( tmp, 2.0f );
    checkCudaErrors( cudaStreamSynchronize( NULL ) );
    test = v.at(0)[0];
    std::cout << " Host TALK: After first kernel launch, value is " << test << std::endl;

    //Initialize second vector on GPU 1, but we do not use it
    cudaSetDevice( 1 );
    v.emplace_back( 65536, 1.0f );
    std::cout << " Host TALK: Raw pointer of vector 0 at step 1 is now " << (void*)thrust::raw_pointer_cast( v.at(0).data() ) << " != " << (void*)tmp << std::endl;
    std::cout << " Host TALK: Raw pointer of vector 1 at step 1 is " << (void*)thrust::raw_pointer_cast( v.at(1).data() ) << std::endl; 

    //Try to use the first vector : No segmentation fault ?
    test = v.at(0)[0];
    std::cout << " Host TALK: Before second kernel launch, value is " << test << std::endl;
    write_memory<<<1,1,0,0>>>( thrust::raw_pointer_cast( v.at(0).data() ), 3.0f );
    checkCudaErrors( cudaStreamSynchronize( NULL ) );
    test = v.at(0)[0];
    std::cout << " Host TALK: After second kernel launch, value is " << test << std::endl;

    //Raw pointer stored elsewhere: generates a segmentation fault
    write_memory<<<1,1,0,0>>>( tmp, 4.0f );
    checkCudaErrors( cudaStreamSynchronize( NULL ) );
    test = v.at(0)[0];
    std::cout << " Host TALK: After third kernel launch, value is " << test << std::endl;

    return 0;
}

这是它在我的机器上产生的输出示例:

主机TALK: vector 0的原始指针在步骤0 0xb043c0000
GPU TALK:原始指针是 0xb043c0000
主机发言:第一次内核启动后，值为 2
主机对话:第 1 步中 vector 0 的原始指针现在为 0xb08000000 != 0xb043c0000
主机TALK:第1步 vector 1的原始指针是0xb07fc0000
主持人发言:在第二个内核启动之前，值为 2
GPU TALK:原始指针为 0xb08000000
主机发言:第二次内核启动后，值为 3
GPU TALK:原始指针是 0xb043c0000
./test.cu(68):CUDA运行时API错误77:遇到非法内存访问抛出“thrust::system::system_error”实例后终止调用what():遇到非法内存访问

提前感谢您的帮助，我也可能会在推力的github上问这个问题。

编辑: 感谢 m.s 和 Hiura，这里的代码可以按预期工作:

/********************************************************************************************
** Compile using nvcc ./test.cu -gencode arch=compute_35,code=sm_35 -std=c++11 -o test.exe **
********************************************************************************************/

//Standard Library
#include <iostream>
#include <vector>

//Cuda
#include "cuda_runtime.h"

//Thrust
#include <thrust/device_vector.h>
#include <thrust/functional.h>
#include <thrust/transform.h>

inline void __checkCudaErrors( cudaError err, const char *file, const int line )
{
    if( err != cudaSuccess )
    {
        printf("%s(%i) : CUDA Runtime API error %i : %s \n",file ,line, (int)err, cudaGetErrorString(err) );
    }
};

#define checkCudaErrors(err)    __checkCudaErrors (err, __FILE__, __LINE__)

__global__ void write_memory( float* buf, float value )
{
    printf("GPU TALK: Raw pointer is %p \n",buf);
    buf[0] = value;
}

int main( int argc, char* argv[] )
{
    //declare a vector of vector
    std::vector<thrust::device_vector<float> > v;
    v.reserve(2);
    float test;
    float* tmp;

    //Initialize first vector on GPU 0
    cudaSetDevice( 0 );
    v.emplace_back( 65536, 1.0f );
    tmp = thrust::raw_pointer_cast( v.at(0).data() );
    std::cout << " Host TALK: Raw pointer of vector 0 at step 0 " << (void*)tmp << std::endl;

    //Try to use it raw pointer
    write_memory<<<1,1,0,0>>>( tmp, 2.0f );
    checkCudaErrors( cudaStreamSynchronize( NULL ) );
    test = v.at(0)[0];
    std::cout << " Host TALK: After first kernel launch, value is " << test << std::endl;

    //Initialize second vector on GPU 1, but we do not use it
    cudaSetDevice( 1 );
    v.emplace_back( 65536, 1.0f );
    std::cout << " Host TALK: Raw pointer of vector 0 at step 1 is now " << (void*)thrust::raw_pointer_cast( v.at(0).data() ) << " != " << (void*)tmp << std::endl;
    std::cout << " Host TALK: Raw pointer of vector 1 at step 1 is " << (void*)thrust::raw_pointer_cast( v.at(1).data() ) << std::endl; 

    //Try to use the first vector : No segmentation fault ?
    cudaSetDevice( 0 );
    test = v.at(0)[0];
    std::cout << " Host TALK: Before second kernel launch, value is " << test << std::endl;
    write_memory<<<1,1,0,0>>>( thrust::raw_pointer_cast( v.at(0).data() ), 3.0f );
    checkCudaErrors( cudaStreamSynchronize( NULL ) );
    test = v.at(0)[0];
    std::cout << " Host TALK: After second kernel launch, value is " << test << std::endl;

    //Raw pointer stored elsewhere: generates a segmentation fault
    write_memory<<<1,1,0,0>>>( tmp, 4.0f );
    checkCudaErrors( cudaStreamSynchronize( NULL ) );
    test = v.at(0)[0];
    std::cout << " Host TALK: After third kernel launch, value is " << test << std::endl;

    return 0;
}

为了简单起见，这是我的代码中最后一个没有使用指向对象的指针 vector 而不是对象 vector 的地方之一，但我发现我应该避免这些烦人的移动/复制问题......

现在的输出是:

主机TALK: vector 0的原始指针在步骤0 0xb043c0000
GPU TALK:原始指针是 0xb043c0000
主机发言:第一次内核启动后，值为 2
主机对话:第 1 步 vector 0 的原始指针现在为 0xb043c0000 != xb043c0000
主机TALK:第1步 vector 1的原始指针是0xb07fc0000
主持人发言:在第二个内核启动之前，值为 2
GPU TALK:原始指针是 0xb043c0000
主机发言:第二次内核启动后，值为 3
GPU TALK:原始指针是 0xb043c0000
主机TALK:第三次内核启动后，值为4

Answer 1

因此我快速安装了 CUDA 来测试我的假设:添加 reserve 语句可以保留地址。

//declare a vector of vector
std::vector<thrust::device_vector<float> > v;
v.reserve(2); // <<-- HERE
float test;
float* tmp;

以及输出，首先没有补丁。

 $ nvcc thrust.cu  -std=c++11 -o test
 $ ./test 
  Host TALK: Raw pointer of vector 0 at step 0 0x700ca0000
 GPU TALK: Raw pointer is 0x700ca0000 
  Host TALK: After first kernel launch, value is 2
  Host TALK: Raw pointer of vector 0 at step 1 is now 0x700d20000 != 0x700ca0000
  Host TALK: Raw pointer of vector 1 at step 1 is 0x700ce0000
  Host TALK: Before second kernel launch, value is 2
 GPU TALK: Raw pointer is 0x700d20000 
  Host TALK: After second kernel launch, value is 3
 GPU TALK: Raw pointer is 0x700ca0000 
  Host TALK: After third kernel launch, value is 3

使用补丁:

 $ nvcc thrust.cu  -std=c++11 -o test
 $ ./test 
  Host TALK: Raw pointer of vector 0 at step 0 0x700ca0000
 GPU TALK: Raw pointer is 0x700ca0000 
  Host TALK: After first kernel launch, value is 2
  Host TALK: Raw pointer of vector 0 at step 1 is now 0x700ca0000 != 0x700ca0000
  Host TALK: Raw pointer of vector 1 at step 1 is 0x700ce0000
  Host TALK: Before second kernel launch, value is 2
 GPU TALK: Raw pointer is 0x700ca0000 
  Host TALK: After second kernel launch, value is 3
 GPU TALK: Raw pointer is 0x700ca0000 
  Host TALK: After third kernel launch, value is 4