CUDA 固定从设备刷新内存
Posted
技术标签:
【中文标题】CUDA 固定从设备刷新内存【英文标题】：CUDA pinned memory flushing from the device 【发布时间】：2013-05-01 07:21:12 【问题描述】：CUDA 5，设备功能 3.5，VS 2012，64 位 Win 2012 服务器。
线程之间没有共享内存访问，每个线程都是独立的。
我正在使用零拷贝的固定内存。只有当我在主机上发出cudaDeviceSynchronize 时，我才能从主机读取设备已写入的固定内存。
我希望能够：
设备更新后立即刷新到固定内存。 不阻塞设备线程（可能通过异步复制）
我尝试在每次设备写入后调用__threadfence_system 和__threadfence，但没有刷新。
以下是演示我的问题的完整示例 CUDA 代码：
#include <conio.h>
#include <cstdio>
#include "cuda.h"
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
__global__ void Kernel(volatile float* hResult)
int tid = threadIdx.x + blockIdx.x * blockDim.x;
printf("Kernel %u: Before Writing in Kernel\n", tid);
hResult[tid] = tid + 1;
__threadfence_system();
// expecting that the data is getting flushed to host here!
printf("Kernel %u: After Writing in Kernel\n", tid);
// time waster for-loop (sleep)
for (int timeWater = 0; timeWater < 100000000; timeWater++);
void main()
size_t blocks = 2;
volatile float* hResult;
cudaHostAlloc((void**)&hResult,blocks*sizeof(float),cudaHostAllocMapped);
Kernel<<<1,blocks>>>(hResult);
int filledElementsCounter = 0;
// naiive thread implementation that can be impelemted using
// another host thread
while (filledElementsCounter < blocks)
// blocks until the value changes, this moves sequentially
// while threads have no order (fine for this sample).
while(hResult[filledElementsCounter] == 0);
printf("%f\n", hResult[filledElementsCounter]);;
filledElementsCounter++;
cudaFreeHost((void *)hResult);
system("pause");
目前此示例将无限期等待，因为除非我发出 cudaDeviceSynchronize，否则不会从设备中读取任何内容。下面的示例有效，但它不是我想要的，因为它违背了异步复制的目的：
void main()
size_t blocks = 2;
volatile float* hResult;
cudaHostAlloc((void**)&hResult, blocks*sizeof(float), cudaHostAllocMapped);
Kernel<<<1,blocks>>>(hResult);
cudaError_t error = cudaDeviceSynchronize();
if (error != cudaSuccess) throw;
for(int i = 0; i < blocks; i++)
printf("%f\n", hResult[i]);
cudaFreeHost((void *)hResult);
system("pause");
【问题讨论】：
这个问题你解决了吗？您是否尝试使用动态并行将数据写入 CPU 主机的内存？在内核函数中使用cudaMemcpyAsync(uva_host_ptr, device_ptr, size);，如下链接所示：on-demand.gputechconf.com/gtc/2012/presentations/…
【参考方案1】：
我在带有 CUDA 5.5 和 Tesla M2090 的 Centos 6.2 上使用您的代码，可以得出以下结论：
它在你的系统上不起作用的问题一定是驱动问题，我建议你获取 TCC 驱动程序。
我附上了我的代码，它运行良好并且可以满足你的需求。这些值在内核结束之前出现在主机端。如您所见，我添加了一些计算代码以防止由于编译器优化而删除 for 循环。我添加了一个流和一个在流中的所有工作完成后执行的回调。程序输出12，很长一段时间内什么都不做，直到stream finished...被打印到控制台。
#include <iostream>
#include "cuda.h"
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#define SEC_CUDA_CALL(val) checkCall ( (val), #val, __FILE__, __LINE__ )
bool checkCall(cudaError_t result, char const* const func, const char *const file, int const line)
if (result != cudaSuccess)
std::cout << "CUDA (runtime api) error: " << func << " failed! " << cudaGetErrorString(result) << " (" << result << ") " << file << ":" << line << std::endl;
return result != cudaSuccess;
class Callback
public:
static void CUDART_CB dispatch(cudaStream_t stream, cudaError_t status, void *userData);
private:
void call();
;
void CUDART_CB Callback::dispatch(cudaStream_t stream, cudaError_t status, void *userData)
Callback* cb = (Callback*) userData;
cb->call();
void Callback::call()
std::cout << "stream finished..." << std::endl;
__global__ void Kernel(volatile float* hResult)
int tid = threadIdx.x + blockIdx.x * blockDim.x;
hResult[tid] = tid + 1;
__threadfence_system();
float A = 0;
for (int timeWater = 0; timeWater < 100000000; timeWater++)
A = sin(cos(log(hResult[0] * hResult[1]))) + A;
A = sqrt(A);
int main(int argc, char* argv[])
size_t blocks = 2;
volatile float* hResult;
SEC_CUDA_CALL(cudaHostAlloc((void**)&hResult,blocks*sizeof(float),cudaHostAllocMapped));
cudaStream_t stream;
SEC_CUDA_CALL(cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking));
Callback obj;
Kernel<<<1,blocks,NULL,stream>>>(hResult);
SEC_CUDA_CALL(cudaStreamAddCallback(stream, Callback::dispatch, &obj, 0));
int filledElementsCounter = 0;
while (filledElementsCounter < blocks)
while(hResult[filledElementsCounter] == 0);
std::cout << hResult[filledElementsCounter] << std::endl;
filledElementsCounter++;
SEC_CUDA_CALL(cudaStreamDestroy(stream));
SEC_CUDA_CALL(cudaFreeHost((void *)hResult));
没有调用返回错误，cuda-memcheck 没有发现任何问题。这按预期工作。你真的应该试试 TCC 驱动程序。
【讨论】：
谢谢！但您的意思可能是Kernel<<<1,threads 而不是Kernel<<<1,blocks？我可以从以下位置下载 TCC 驱动程序：nvidia.com/object/software-for-tesla-products.html 但是我可以将它用于 nVidia Quadro Mobile 还是我必须使用 GPU nVidia Quadro (Kepler GK107/GK106) 来解决这个问题？
代码是从原始问题中复制的，但是第二个参数是用于线程的。我对 TCC 驱动程序没有经验，但我认为它也适用于 Quadros。看这里：***.com/questions/19098650/…【参考方案2】：
您不能将主机指针直接传递给内核。如果您使用带有cudaHostAllocMapped 标志的cudaHostAlloc 分配主机内存，那么首先您必须检索映射的主机内存的设备指针，然后才能在内核中使用它。使用cudaHostGetDevicePointer获取映射主机内存的设备指针。
float* hResult, *dResult;
cudaHostAlloc((void**)&hResult, blocks*sizeof(float), cudaHostAllocMapped);
cudaHostGetDevicePointer(&dResult,hResult);
Kernel<<<1,blocks>>>(dResult);
【讨论】：
当你说“你不能通过”时，你的意思是解决我的冲洗问题还是一般来说？因为当我用cudaDeviceSynchronize 替换我的while 循环时，我可以在不执行任何cudaMemcpy 的情况下访问hResult 中的数据。我仍然看不到您建议的解决方案如何解决冲洗问题。我是否会一直在 dResult 上执行 cudaMemcpyAsync，直到我在其中找到一些东西？
实际上我指出了一个一般错误，它会导致未定义的行为。刷新问题可能是由于内核中的 printf 语句造成的。因为一旦内核完成执行，内核中的printf 就会转储其输出。
我在问题中添加了另一个示例，该示例有效，但同步。您是在告诉我第二个样本的行为未定义吗？它正在工作，即使我删除了内核 printf
只有当您没有有统一的虚拟地址时才适用。如果您有 UVA（即 sm_20 或更高版本、64 位 Linux 或带有 TCC/WinXP 的 64 位 Windows），则无需致电 cudaHostGetDevicePointer()。请参阅docs.nvidia.com/cuda/cuda-c-programming-guide/… 了解更多信息。
@Tom，所以除了 TCC，我拥有所有这些，因为我目前正在使用 GTX Titan，但代码仍然可以正常工作，而无需调用 cudaHostGetDevicePointer，这是侥幸吗？【参考方案3】：
调用__threadfence_system() 将确保在继续之前写入对系统可见，但您的CPU 将缓存h_result 变量，因此您只是在无限循环中旋转旧值。尝试将 h_result 标记为 volatile。
【讨论】：
我已经更新了上面的示例并添加了 __threadfence_system() 和 volatile，因为添加 volatile 是个好主意。但是，我仍然无法阅读任何内容。以上是关于CUDA 固定从设备刷新内存的主要内容，如果未能解决你的问题，请参考以下文章