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