推力::max_element 是在主机还是设备上运行？答案

【问题标题】：Is thrust::max_element run on host or device?推力::max_element 是在主机还是设备上运行？
【发布时间】：2016-01-29 14:14:17
【问题描述】：

我需要在存储在我的设备上的长数组中找到最大元素。我想我可以使用推力::max_element 来做到这一点。我在下面代码的 while 循环中调用了thrust::max_element。我只是给它两个设备指针（注意 real 只是 float 的 typedef）。我能不能只传递推力::max_element 设备指针。它是否试图在主机上找到最大元素？我问这个是因为我的代码在那之后因为段错误而失败。

int main()
{
    cuda_error(cudaSetDevice(1), "set device");
    const size_t DIM = 50;

    real* grid_d;
    real* next_grid_d;
    cuda_error(cudaMalloc(&grid_d, sizeof(real) * DIM * DIM * DIM), "malloc grid");
    cuda_error(cudaMalloc(&next_grid_d, sizeof(real) * DIM * DIM * DIM), "malloc next grid");
    cuda_error(cudaMemset(grid_d, 0, sizeof(real) * DIM * DIM * DIM), "memset grid");

    ConstantSum point_charge(0.3, DIM / 2, DIM / 2, DIM / 2);
    ConstantSum* point_charge_d;
    cuda_error(cudaMalloc(&point_charge_d, sizeof(ConstantSum)), "malloc constant sum");
    cuda_error(cudaMemcpy(point_charge_d, &point_charge, sizeof(ConstantSum), cudaMemcpyHostToDevice), "memset constant sum");

    real max_err;
    do
    {
        compute_next_grid_kernel<<< DIM, dim3(16, 16) >>>(grid_d, next_grid_d, DIM, point_charge_d, 1);
        cuda_error(cudaGetLastError(), "kernel launch");
        max_err = *thrust::max_element(grid_d, grid_d + DIM * DIM * DIM);

        std::swap(grid_d, next_grid_d);
    }
    while(max_err > 0.1);

    real* frame = new real[DIM * DIM];
    cuda_error(cudaMemcpy(frame, grid_d + DIM * DIM * (DIM / 2), DIM * DIM * sizeof(real), cudaMemcpyDeviceToHost), "memcpy frame");

    cuda_error(cudaFree(grid_d), "free grid");
    cuda_error(cudaFree(next_grid_d), "free next grid");
    cuda_error(cudaFree(point_charge_d), "free point charge");

    for(int i = 0; i < DIM; i++)
    {
        for(int j = 0; j < DIM; j++)
        {
            std::cout << frame[DIM * i + j] << "\t";
        }
        std::cout << "\n";
    }

    delete[] frame;
    return 0;
}

【问题讨论】：

标签： c++ cuda thrust

【解决方案1】：

一般情况下，thrust 使用传递的迭代器的类型来确定算法后端将在主机或设备上运行（在最新版本中也有无标签和显式执行策略选择，但这是一个不同的讨论)。

在你的情况下，因为 grid_d 是一个 host 指针（它的 value 是主机地址还是设备地址无关紧要），thrust 将尝试运行算法在主机上。这是段错误的来源，您正在尝试访问主机上的设备地址。

要完成这项工作，您需要将指针转换为 thrust::dev_ptr，类似于：

thrust::device_ptr<real> grid_start = thrust::device_pointer_cast(grid_d);
thrust::device_ptr<real> grid_end= thrust::device_pointer_cast(grid_d + DIM * DIM * DIM);

auto max_it = thrust::max_element(grid_start, grid_end);
max_error = *max_it;

[警告，在浏览器中编写，从未编译或测试，使用风险自负]

通过传递thrust::dev_ptr，发生正确的标签选择，并且闭包将在设备上运行。

另一种不进行强制转换的解决方案是指定execution policy device：

thrust::max_element(thrust::device, grid_d, grid_d + DIM * DIM * DIM);

没有，只有 Thrust 1.7 及更高版本才支持显式执行策略控制。

【讨论】：