c++ - thrust::max_element 是在主机上运行还是在设备上运行？

Question

我需要找到存储在我设备上的长数组中的最大元素。我想我可以使用 thrust::max_element 来做到这一点。我在下面代码的 while 循环中调用 thrust::max_element。我只是给它两个设备指针(注意 real 只是 float 的 typedef)。我可以不只传递 thrust::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;
}

Answer 1

一般来说，thrust是通过传递的迭代器的类型来决定算法后端是运行在host上还是device上(最新版本也有tag free和explicit execution policy selection，但那是另外的讨论).

在你的例子中，因为 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 及更高版本才支持显式执行策略控制。