thrust :: raw_pointer_cast和多GPU，奇怪的行为

Question

我在代码中使用了很多东西，因为它是一个很好的包装器并提供了非常有用的实用程序，因为支持异步行为已被添加，所以我更加确信。

我的代码运行良好，使用cuda推力直到我最近试图在我的应用程序中添加多GPU支持。 我经历恼人的

CUDA运行时API错误77：一个非法的内存访问时遇到

在我的代码，以前从来没有表现出任何边界问题的一部分。

我在我的代码中添加了冗长的内容，看起来我的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：矢量0的原始指针在步骤0 0xb043c0000
GPU TALK：原始指针是0xb043c0000
主机TALK：经过第一内核启动，值是2
主机TALK：在步骤1矢量0的原始指针现在是0xb08000000 = 0xb043c0000
主机TALK：在步骤1载体1的原始指针是0xb07fc0000
主机TALK：第二内核启动之前，值为2
GPU TALK ：原始指针是0xb08000000
主机快讯：第二内核启动后，值3
GPU快讯：原始指针是0xb043c0000
./test.cu(68）：CUDA运行时API错误77：遇到一个非法的内存访问在抛出'thrust :: system :: system_error'实例后终止调用what（）：遇到非法内存访问

非常感谢您的帮助，我也可以在推特的github上提出这个问题。

编辑： 由于MS和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; 
} 

这是我的代码的最后一个地方，我没有用指针的矢量对象，而不是一个简单的对象矢量，但我看到，我应该不得不避免这些恼人的移动/复制问题...

输出现在是：

主机TALK：在步骤矢量0的原始指针0 0xb043c0000
GPU TALK：原始指针是0xb043c0000
主机TALK：第一内核启动后，值为2
主机TALK ：在步骤1矢量0的原始指针现在是0xb043c0000 = xb043c0000
主机TALK：在步骤1载体1的原始指针是0xb07fc0000
主机TALK：第二内核启动之前，值为2
GPU TALK：原始指针是0xb043c0000
主机快讯：第二内核启动后，值3
GPU快讯：原始指针是0xb043c0000
主机快讯：第三内核启动后，值为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