CUDA学习（九十八）

流关联示例：
将数据与流关联可以对CPU + GPU并发性进行细粒度的控制，但是在使用低于6.x的计算能力的设备时，必须记住哪些数据是可见的。 查看较早的同步示例：

__device__ __managed__ int x, y = 2;
__global__ void kernel() {
    x = 10;
}
int main() {
    cudaStream_t stream1;
    cudaStreamCreate(&stream1);
    cudaStreamAttachMemAsync(stream1, &y, 0, cudaMemAttachHost);
    cudaDeviceSynchronize(); // Wait for Host attachment to occur.
    kernel << < 1, 1, 0, stream1 >> >(); // Note: Launches into stream1.
    y = 20; // Success – a kernel is running but “y”
            // has been associated with no stream.
    return 0;
}

在这里，我们明确地将y与主机可访问性相关联，从而使CPU始终能够访问。 （如前所述，请注意在访问之前没有cudaDeviceSynchronize（）。）运行内核的GPU访问y现在会产生未定义的结果。
请注意，将变量与流关联不会更改任何其他变量的关联。 例如。 将x关联到stream1并不能确保只有x被stream1中启动的内核访问，因此错误是由以下代码引起的：

__device__ __managed__ int x, y = 2;
__global__ void kernel() {
    x = 10;
}
int main() {
    cudaStream_t stream1;
    cudaStreamCreate(&stream1);
    cudaStreamAttachMemAsync(stream1, &x);// Associate “x” with stream1.
    cudaDeviceSynchronize(); // Wait for “x” attachment to occur.
    kernel << < 1, 1, 0, stream1 >> >(); // Note: Launches into stream1.
    y = 20; // ERROR: “y” is still associated
    globally
        // with all streams by default
        return 0;
}

请注意，对y的访问将导致错误，因为尽管x已经与流相关联，但我们已经告诉系统没有人可以看到y。 因此，系统保守地认为内核可能会访问它并阻止CPU这样做。
流连接与多线程主机程序：
cudaStreamAttachMemAsync（）的主要用途是使用CPU线程启用独立任务并行。 通常在这样的程序中，CPU线程为其生成的所有工作创建自己的流，因为使用CUDA的NULL流将导致线程之间的依赖关系。
受管理数据对任何GPU流的默认全局可见性可能会使得难以避免多线程程序中的CPU线程之间的交互。 函数cudaStreamAttachMemAsync（）因此用于将线程的托管分配与该线程自己的流相关联，并且该关联在线程的整个生命周期中通常不会更改。
这样的程序只需将一个调用添加到cudaStreamAttachMemAsync（）中即可为其数据访问使用统一内存：

// This function performs some task, in its own private stream.
void run_task(int *in, int *out, int length) {
    // Create a stream for us to use.
    cudaStream_t stream;
    cudaStreamCreate(&stream);
    // Allocate some managed data and associate with our stream.
    // Note the use of the host-attach flag to cudaMallocManaged();
    // we then associate the allocation with our stream so that
    // our GPU kernel launches can access it.
    int *data;
    cudaMallocManaged((void **)&data, length, cudaMemAttachHost);
    cudaStreamAttachMemAsync(stream, data);
    cudaStreamSynchronize(stream);
    // Iterate on the data in some way, using both Host & Device.
    for (int i = 0; i<N; i++) {
        transform << < 100, 256, 0, stream >> >(in, data, length);
        cudaStreamSynchronize(stream);
        host_process(data, length); // CPU uses managed data.
        convert << < 100, 256, 0, stream >> >(out, data, length);
    }
    cudaStreamSynchronize(stream);
    cudaStreamDestroy(stream);
    cudaFree(data);
}

在这个例子中，分配流关联只建立一次，然后数据被主机和设备重复使用。 结果是比在主机和设备之间明确复制数据时更简单的代码，尽管结果是相同的。