c++ - 在 CUDA 6 中结合纹理内存 Unified Memory

Question

我正在使用 CUDA 6 为 Jetson TK1 编写 CUDA 应用程序。我从 Mark Harris 的博客文章中得到了印象

Jetson TK1: Mobile Embedded Supercomputer Takes CUDA Everywhere

Tegra K1 的内存在物理上是统一的。我还观察到结果表明 cudaMallocManaged 对于全局内存比普通 cudaMemcpy 快得多。这可能是因为统一内存不需要任何复制。

但是，当我想为我的部分应用程序使用纹理内存时，我该怎么办？我没有发现任何使用 cudaMallocManaged 的纹理支持，所以我假设我必须使用普通的 cudaMemcpyToArray 和 bindTextureToArray？

使用前面提到的方法似乎经常有效，但由 cudaMallocManaged 管理的变量有时会给我带来奇怪的段错误。这是将纹理内存与统一内存一起使用的正确方法吗？下面的代码说明了我是如何做到的。这段代码工作正常，但我的问题是这是否是正确的方法，或者它是否可能会产生未定义的行为，从而导致例如段错误。

#define width 16
#define height 16
texture<float, cudaTextureType2D, cudaReadModeElementType> input_tex;

__global__ void some_tex_kernel(float* output){
    int i= threadIdx.x;
    float x = i%width+0.5f;
    float y =  i/width+0.5f;
    output[i] = tex2D(input_tex, x, y);
}

int main(){
    float* out;
    if(cudaMallocManaged(&out, width*height*sizeof(float))!= cudaSuccess)
        std::cout << "unified not working\n";

    for(int i=0; i< width*height; ++i){
        out[i] = float(i);
    }

    const cudaChannelFormatDesc desc = cudaCreateChannelDesc<float>();
    cudaArray* input_t;
    cudaMallocArray(&input_t, &desc, width, height);
    cudaMemcpyToArrayAsync(input_t, 0, 0, out, width*height*sizeof(float),  cudaMemcpyHostToDevice);

    input_tex.filterMode = cudaFilterModeLinear;
    cudaBindTextureToArray(input_tex, input_t, desc);

    some_tex_kernel<<<1, width*height>>>(out);
    cudaDeviceSynchronize();

    for(int i=0;i<width*height; ++i)
        std::cout << out[i] << " ";

    cudaFree(out);
    cudaFreeArray(input_t); 
    }
}

我觉得奇怪的另一件事是，如果我删除代码中的 cudaDeviceSynchronize()，我总是会遇到段错误。我知道如果我在没有同步的情况下读取结果可能无法完成，但变量不应该仍然可以访问吗？

有人知道吗？

马蒂亚斯

Answer 1

only managed memory possibilities at this time是使用 __device__ __managed__ 的静态分配或使用 cudaMallocManaged() 的动态分配。不直接支持纹理、表面、常量内存等。

您对纹理的使用很好。纹理使用和托管内存之间的唯一重叠是在以下调用中:

cudaMemcpyToArrayAsync(input_t, 0, 0, out, width*height*sizeof(float),  cudaMemcpyHostToDevice);

其中托管内存是传输的来源(即主机端)。只要在没有内核执行的期间发出调用，这是可以接受的(见下文)。

"Another thing that I find odd is that if I remove the cudaDeviceSynchronize() in the code I always get segmentation faults."

cudaDeviceSynchronize(); 在内核调用之后是必需的，以使托管内存再次对主机可见。我建议你阅读 this section of the documentation小心:

"In general, it is not permitted for the CPU to access any managed allocations or variables while the GPU is active. Concurrent CPU/GPU accesses, ... will cause a segmentation fault..."

正如您所指出的，您发布的代码工作正常。如果您有其他代码在使用托管内存时出现不可预测的段错误，我会仔细检查代码流(特别是如果您使用流，即并发)以确保主机仅在 cudaDeviceSynchronize( ); 已发出，并且在任何后续内核调用之前。