thrust::transform() 导致从主机到设备的 cudaErrorIllegalAddress

Question

以下

test.cu

程序

#include <thrust/copy.h>
#include <thrust/execution_policy.h>
#include <thrust/transform.h>
#include <thrust/device_vector.h>
#include <thrust/host_vector.h>
#include <iostream>

using HOST_TYPE=int32_t;
using DEVICE_TYPE=int;

template <typename T>
struct Cast {
  __host__ __device__ T operator()(HOST_TYPE i) const {
    return static_cast<T>(i);
  }
};

int main() {
    // Initialize host data
    thrust::host_vector<HOST_TYPE> const h_vec{1, 2, 3, 4, 5};
    
    // Allocate space on the device
    thrust::device_vector<DEVICE_TYPE> device_data(h_vec.size());

    // Copy data from host to device
    //thrust::copy(h_vec.cbegin(), h_vec.cend(), device_data.begin());  // this works
    thrust::transform(h_vec.cbegin(), h_vec.cend(), device_data.begin(), Cast<DEVICE_TYPE>{});
    
    // Copy back to host to check
    thrust::host_vector<DEVICE_TYPE> host_data_copy = device_data;
    for (DEVICE_TYPE val : host_data_copy) {
        std::cout << val << " ";
    }
    std::cout << std::endl;
    
    return 0;
}

原因

$ nvcc test.cu
$ ./a.out 
terminate called after throwing an instance of 'thrust::system::system_error'
  what():  parallel_for: failed to synchronize: cudaErrorIllegalAddress: an illegal memory access was encountered
Aborted (core dumped)

这发生在生产线上

thrust::transform(h_vec.cbegin(), h_vec.cend(), device_data.begin(), Cast<DEVICE_TYPE>{});

即使类似的

thrust::copy()

运行良好：

thrust::copy(h_vec.cbegin(), h_vec.cend(), device_data.begin());  // this works

我在文档中找不到任何内容说明

thrust::transform()

不应在设备和主机之间转换数据。我是不是在某个地方错过了这个？

使用

thrust::host

或

thrust::device

执行策略没有帮助。

版本：

$ nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2023 NVIDIA Corporation
Built on Fri_Nov__3_17:16:49_PDT_2023
Cuda compilation tools, release 12.3, V12.3.103
Build cuda_12.3.r12.3/compiler.33492891_0

注意：实际应用需要

HOST_TYPE=char

，但为了调试/说明目的，将其更改为

HOST_TYPE=int32_t

，并与

std::copy()

进行比较。

Answer 1

请参阅转换：

除了可以在主机和设备之间复制数据的 Thrust::copy 之外，Thrust 算法的所有迭代器参数都应该位于同一位置：要么全部位于主机上，要么全部位于设备上。当违反此要求时，编译器将产生错误消息。