thrust::raw_pointer_cast 和多个 GPU，奇怪的行为
thrust::raw_pointer_cast and multiple GPU, weird behaviour
我在我的代码中经常使用 thrust，因为它是一个很棒的包装器并提供了非常有用的实用程序，因为添加了对异步行为的支持，我更加确信。
在我最近尝试在我的应用程序中添加多 GPU 支持之前，我的代码使用 cuda 推力运行良好。
我经历了烦人的事
CUDA Runtime API error 77 : an illegal memory access was encountered
我的部分代码以前从未出现过任何边界问题。
我在我的代码中添加了冗长的代码，看起来我的 thrust::device_vector 指针地址在执行过程中发生了变化，没有明显的原因，在手写内核中生成了错误 77。
我可能误解了 UVA 概念及其最终结果 "side effects"，但是，我仍然有兴趣了解导致指针更新的过程。
我无法准确重现我的问题，其中我没有使用临时主机变量来存储 cuda 内存指针，而是仅在内核包装器调用需要时即时 thrust::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;
}
这是它在我的机器上产生的输出示例：
主机对话：第 0 步向量 0 的原始指针 0xb043c0000
GPU 对话：原始指针是 0xb043c0000
Host TALK：第一次内核启动后，值为 2
主机对话：第 1 步向量 0 的原始指针现在是 0xb08000000 != 0xb043c0000
主机谈话：第 1 步向量 1 的原始指针是 0xb07fc0000
Host TALK：在第二次内核启动之前，值为 2
GPU 对话：原始指针是 0xb08000000
Host TALK：第二次内核启动后，值为 3
GPU 对话：原始指针是 0xb043c0000
./test.cu(68)：CUDA 运行时 API 错误 77：遇到非法内存访问 在抛出 'thrust::system::system_error' what() 的实例后终止调用：遇到非法内存访问
在此先感谢您的帮助，我也可能会在 thrust 上问这个问题 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;
}
这是我代码中最后一个地方，为了简单起见，我没有使用指向对象的指针向量而不是对象向量，但我发现我应该避免这些恼人的 move/copy 问题...
现在的输出是：
主机对话：第 0 步向量 0 的原始指针 0xb043c0000
GPU 对话：原始指针是 0xb043c0000
Host TALK：第一次内核启动后，值为 2
主机对话：第 1 步向量 0 的原始指针现在是 0xb043c0000 != xb043c0000
主机谈话：第 1 步向量 1 的原始指针是 0xb07fc0000
Host TALK：在第二次内核启动之前，值为 2
GPU 对话：原始指针是 0xb043c0000
Host TALK：第二次内核启动后，值为 3
GPU 对话：原始指针是 0xb043c0000
主持人谈话：第三次内核启动后，值为 4
所以我快速安装了 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
我在我的代码中经常使用 thrust，因为它是一个很棒的包装器并提供了非常有用的实用程序，因为添加了对异步行为的支持，我更加确信。
在我最近尝试在我的应用程序中添加多 GPU 支持之前，我的代码使用 cuda 推力运行良好。 我经历了烦人的事
CUDA Runtime API error 77 : an illegal memory access was encountered
我的部分代码以前从未出现过任何边界问题。
我在我的代码中添加了冗长的代码，看起来我的 thrust::device_vector 指针地址在执行过程中发生了变化，没有明显的原因，在手写内核中生成了错误 77。
我可能误解了 UVA 概念及其最终结果 "side effects"，但是，我仍然有兴趣了解导致指针更新的过程。
我无法准确重现我的问题，其中我没有使用临时主机变量来存储 cuda 内存指针，而是仅在内核包装器调用需要时即时 thrust::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;
}
这是它在我的机器上产生的输出示例：
主机对话：第 0 步向量 0 的原始指针 0xb043c0000
GPU 对话：原始指针是 0xb043c0000
Host TALK：第一次内核启动后，值为 2
主机对话：第 1 步向量 0 的原始指针现在是 0xb08000000 != 0xb043c0000
主机谈话：第 1 步向量 1 的原始指针是 0xb07fc0000
Host TALK：在第二次内核启动之前，值为 2
GPU 对话：原始指针是 0xb08000000
Host TALK：第二次内核启动后，值为 3
GPU 对话：原始指针是 0xb043c0000
./test.cu(68)：CUDA 运行时 API 错误 77：遇到非法内存访问 在抛出 'thrust::system::system_error' what() 的实例后终止调用：遇到非法内存访问
在此先感谢您的帮助，我也可能会在 thrust 上问这个问题 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;
}
这是我代码中最后一个地方，为了简单起见，我没有使用指向对象的指针向量而不是对象向量，但我发现我应该避免这些恼人的 move/copy 问题...
现在的输出是：
主机对话：第 0 步向量 0 的原始指针 0xb043c0000
GPU 对话：原始指针是 0xb043c0000
Host TALK：第一次内核启动后，值为 2
主机对话：第 1 步向量 0 的原始指针现在是 0xb043c0000 != xb043c0000
主机谈话：第 1 步向量 1 的原始指针是 0xb07fc0000
Host TALK：在第二次内核启动之前，值为 2
GPU 对话：原始指针是 0xb043c0000
Host TALK：第二次内核启动后，值为 3
GPU 对话：原始指针是 0xb043c0000
主持人谈话：第三次内核启动后，值为 4
所以我快速安装了 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