CUDA固定内存从设备中刷新

Question

CUDA 5，设备功能3.5，VS 2012,64位Win 2012 Server。

线程之间没有共享内存访问，每个线程都是独立的。

我使用零拷贝的固定内存。在主机上，只有当我在主机上发出cudaDeviceSynchronize时，我才能读取设备写入的固定内存。

我希望能够到：

1. 水冲到锁定的存储，一旦设备已经更新了它。
2. 不会阻止设备线程（可能由异步复制）

我打过电话__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"); 
} 

Answer 1

您不能直接通过主机指针到内核。如果使用cudaHostAlloc和cudaHostAllocMapped标志分配主机内存，则首先必须检索映射主机内存的设备指针，然后才能在内核中使用它。使用cudaHostGetDevicePointer获取映射主机内存的设备指针。

float* hResult, *dResult; 
cudaHostAlloc((void**)&hResult, blocks*sizeof(float), cudaHostAllocMapped); 
cudaHostGetDevicePointer(&dResult,hResult); 
Kernel<<<1,blocks>>>(dResult); 

Answer 2

调用__threadfence_system()将确保写入对系统可见继续之前，但你的CPU会被缓存h_result变量，因此你只是在一个无限循环旋转的旧值。尝试将h_result标记为volatile。

Answer 3

我用你的代码上播放一个CentOS 6.2 CUDA 5.5和特斯拉M2090，可以断定这一点：

它不会在您的系统的问题必须是驱动的问题，我建议你吃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驱动程序。