2

CUDA 5，设备功能 3.5，VS 2012，64 位 Win 2012 服务器。

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

我正在使用零拷贝的固定内存。只有当我在主机上发出 a 时，我才能从主机读取设备写入的固定内存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");
}
4

3 回答 3

4

我在带有 CUDA 5.5 和 Tesla M2090 的 Centos 6.2 上使用了您的代码，可以得出以下结论：

它在您的系统上不起作用的问题必须是驱动程序问题，我建议您获取 TCC 驱动程序。

我附上了我的代码，它运行良好，可以做你想做的事。这些值在内核结束之前出现在主机端。如您所见，我添加了一些计算代码以防止由于编译器优化而删除 for 循环。我添加了一个流和一个在流中的所有工作完成后执行的回调。程序输出1 2并在很长一段时间内什么都不做，直到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 驱动程序。

于 2013-10-15T12:27:10.637 回答
2

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

于 2013-05-07T11:58:35.923 回答
2

您不能将主机指针直接传递给内核。如果使用cudaHostAllocwithcudaHostAllocMapped标志分配主机内存，则首先必须检索映射的主机内存的设备指针，然后才能在内核中使用它。用于cudaHostGetDevicePointer获取映射主机内存的设备指针。

float* hResult, *dResult;
cudaHostAlloc((void**)&hResult, blocks*sizeof(float), cudaHostAllocMapped);
cudaHostGetDevicePointer(&dResult,hResult);
Kernel<<<1,blocks>>>(dResult);
于 2013-05-07T11:16:04.133 回答