为什么推力::device_vector 似乎没有机会保存指向其他 device_vector 的原始指针？

Question

我有一个问题，我在其中找到了很多线程，但没有一个明确回答我的问题。 我正在尝试使用推力在 GPU 内核中创建一个多维数组。展平会很困难，因为所有尺寸都是不均匀的，我会上升到 4D。现在我知道我不能拥有 device_vectors 的 device_vectors，无论出于何种根本原因（欢迎解释），所以我尝试绕过原始指针。

我的理由是，原始指针指向 GPU 上的内存，否则我为什么能够从内核中访问它。所以从技术上讲，我应该能够拥有一个 device_vector，它包含原始指针，所有指针都应该可以从 GPU 中访问。这样我构造了以下代码：

thrust::device_vector<Vector3r*> d_fluidmodelParticlePositions(nModels);
thrust::device_vector<unsigned int***> d_allFluidNeighborParticles(nModels);
thrust::device_vector<unsigned int**> d_nFluidNeighborsCrossFluids(nModels);

for(unsigned int fluidModelIndex = 0; fluidModelIndex < nModels; fluidModelIndex++)
{
    FluidModel *model = sim->getFluidModelFromPointSet(fluidModelIndex);
    const unsigned int numParticles = model->numActiveParticles();

    thrust::device_vector<Vector3r> d_neighborPositions(model->getPositions().begin(), model->getPositions().end());
    d_fluidmodelParticlePositions[fluidModelIndex] = CudaHelper::GetPointer(d_neighborPositions);

    thrust::device_vector<unsigned int**> d_fluidNeighborIndexes(nModels);
    thrust::device_vector<unsigned int*> d_nNeighborsFluid(nModels);

    for(unsigned int pid = 0; pid < nModels; pid++)
    {
        FluidModel *fm_neighbor = sim->getFluidModelFromPointSet(pid);

        thrust::device_vector<unsigned int> d_nNeighbors(numParticles);
        thrust::device_vector<unsigned int*> d_neighborIndexesArray(numParticles);

        for(unsigned int i = 0; i < numParticles; i++)
        {
            const unsigned int nNeighbors = sim->numberOfNeighbors(fluidModelIndex, pid, i);        
            d_nNeighbors[i] = nNeighbors;

            thrust::device_vector<unsigned int> d_neighborIndexes(nNeighbors);

            for(unsigned int j = 0; j < nNeighbors; j++)
            {
                d_neighborIndexes[j] = sim->getNeighbor(fluidModelIndex, pid, i, j);
            }

            d_neighborIndexesArray[i] = CudaHelper::GetPointer(d_neighborIndexes);
        }

        d_fluidNeighborIndexes[pid] = CudaHelper::GetPointer(d_neighborIndexesArray);
        d_nNeighborsFluid[pid] = CudaHelper::GetPointer(d_nNeighbors);
    }

    d_allFluidNeighborParticles[fluidModelIndex] = CudaHelper::GetPointer(d_fluidNeighborIndexes);
    d_nFluidNeighborsCrossFluids[fluidModelIndex] = CudaHelper::GetPointer(d_nNeighborsFluid);
}

现在编译器不会抱怨了，但是从内核中访问例如 d_nFluidNeighborsCrossFluids 会起作用，但会返回错误的值。我像这样访问它（同样，从内核中）：

d_nFluidNeighborsCrossFluids[iterator1][iterator2][iterator3];
// Note: out of bounds indexing guaranteed to not happen, indexing is definitely right

问题是，为什么它返回错误的值？我认为它背后的逻辑应该有效，因为我的索引是正确的，并且指针应该是内核中的有效地址。

感谢您抽出宝贵的时间，祝您有美好的一天。

编辑： 这是一个最小的可重现示例。出于某种原因，尽管与我的代码具有相同的结构，但这些值看起来是正确的，但是 cuda-memcheck 揭示了一些错误。取消注释两条注释行会导致我试图解决我的主要问题。这里的 cuda-memcheck 告诉我什么？

/* Part of this example has been taken from code of Robert Crovella 
   in a comment below */
#include <thrust/device_vector.h>
#include <stdio.h>

template<typename T>
static T* GetPointer(thrust::device_vector<T> &vector)
{
  return thrust::raw_pointer_cast(vector.data());
}

__global__ 
void k(unsigned int ***nFluidNeighborsCrossFluids, unsigned int ****allFluidNeighborParticles){

  const unsigned int i = blockIdx.x*blockDim.x + threadIdx.x;

  if(i > 49)
    return;

  printf("i: %d nNeighbors: %d\n", i, nFluidNeighborsCrossFluids[0][0][i]);

  //for(int j = 0; j < nFluidNeighborsCrossFluids[0][0][i]; j++)
  //  printf("i: %d j: %d neighbors: %d\n", i, j, allFluidNeighborParticles[0][0][i][j]);
}


int main(){

  const unsigned int nModels = 2;
  const int numParticles = 50;

  thrust::device_vector<unsigned int**> d_nFluidNeighborsCrossFluids(nModels);
  thrust::device_vector<unsigned int***> d_allFluidNeighborParticles(nModels);

  for(unsigned int fluidModelIndex = 0; fluidModelIndex < nModels; fluidModelIndex++)
  {
    thrust::device_vector<unsigned int*> d_nNeighborsFluid(nModels);
    thrust::device_vector<unsigned int**> d_fluidNeighborIndexes(nModels);

    for(unsigned int pid = 0; pid < nModels; pid++)
    {

      thrust::device_vector<unsigned int> d_nNeighbors(numParticles);
      thrust::device_vector<unsigned int*> d_neighborIndexesArray(numParticles);

      for(unsigned int i = 0; i < numParticles; i++)
      {
        const unsigned int nNeighbors = i;        
        d_nNeighbors[i] = nNeighbors;

        thrust::device_vector<unsigned int> d_neighborIndexes(nNeighbors);

                for(unsigned int j = 0; j < nNeighbors; j++)
                {
                    d_neighborIndexes[j] = i + j;
        }
        d_neighborIndexesArray[i] = GetPointer(d_neighborIndexes);
      }
      d_nNeighborsFluid[pid] = GetPointer(d_nNeighbors);
      d_fluidNeighborIndexes[pid] = GetPointer(d_neighborIndexesArray);
    }
    d_nFluidNeighborsCrossFluids[fluidModelIndex] = GetPointer(d_nNeighborsFluid);
    d_allFluidNeighborParticles[fluidModelIndex] = GetPointer(d_fluidNeighborIndexes);

  }

  k<<<256, 256>>>(GetPointer(d_nFluidNeighborsCrossFluids), GetPointer(d_allFluidNeighborParticles));

  if (cudaGetLastError() != cudaSuccess) 
    printf("Sync kernel error: %s\n", cudaGetErrorString(cudaGetLastError()));

  cudaDeviceSynchronize();
}

Answer 1

你真的应该提供一个minimal, complete, verifiable/reproducible example;你的既不是最小的，也不是完整的，也不是可验证的。

不过，我会回答你的附带问题：

我知道我不能拥有device_vectors 中的device_vectors，无论出于何种根本原因（欢迎解释）

虽然 device_vector 涉及 GPU 上的一堆数据，但它是主机端数据结构 - 否则您将无法在主机端代码中使用它。在主机端，它所拥有的应该是：容量、元素大小、指向实际数据的设备端指针，也许还有更多信息。这类似于std::vector 变量如何引用堆上的数据，但是如果您在本地创建变量，我上面提到的字段将存在于堆栈中。

现在，位于主机内存中的设备向量字段通常无法从设备端访问。在设备端代码中，您通常会使用指向device_vector 管理的设备端数据的原始指针。

另外，请注意，如果您有 thrust::device_vector<T> v，则每次使用 operator[] 都意味着一组单独的 CUDA 调用来将数据复制到设备或从设备复制数据（除非在后台进行一些缓存）。所以你真的想避免在这种结构中使用方括号。

最后，请记住，指针追踪可能会成为性能杀手，尤其是在 GPU 上。您可能需要考虑对您的数据结构进行一些调整，以使其易于扁平化。

Answer 2

device_vector 是一个类定义。该类具有与之关联的各种方法和运算符。允许你这样做的东西：

d_nFluidNeighborsCrossFluids[...]...;

是方括号运算符。该操作员是主机操作员（仅限）。它在设备代码中不可用。此类问题会导致“thrust::device_vector 在设备代码中不可用”的一般性陈述。 device_vector object 本身通常不可用。但是，如果您尝试通过原始指针访问它，则它包含的数据可在设备代码中使用。

这是一个推力设备向量的示例，其中包含指向其他设备向量中包含的数据的指针数组。只要您不尝试使用推力::device_vector 对象本身，该数据就可以在设备代码中使用：

$ cat t1509.cu
#include <thrust/device_vector.h>
#include <stdio.h>

template <typename T>
__global__ void k(T **data){

  printf("the first element of vector 1 is: %d\n", (int)(data[0][0]));
  printf("the first element of vector 2 is: %d\n", (int)(data[1][0]));
  printf("the first element of vector 3 is: %d\n", (int)(data[2][0]));
}


int main(){

  thrust::device_vector<int> vector_1(1,1);
  thrust::device_vector<int> vector_2(1,2);
  thrust::device_vector<int> vector_3(1,3);

  thrust::device_vector<int *> pointer_vector(3);
  pointer_vector[0] = thrust::raw_pointer_cast(vector_1.data());
  pointer_vector[1] = thrust::raw_pointer_cast(vector_2.data());
  pointer_vector[2] = thrust::raw_pointer_cast(vector_3.data());

  k<<<1,1>>>(thrust::raw_pointer_cast(pointer_vector.data()));
  cudaDeviceSynchronize();
}

$ nvcc -o t1509 t1509.cu
$ cuda-memcheck ./t1509
========= CUDA-MEMCHECK
the first element of vector 1 is: 1
the first element of vector 2 is: 2
the first element of vector 3 is: 3
========= ERROR SUMMARY: 0 errors
$

编辑：在您现在发布的 mcve 中，您指出代码的普通运行似乎给出了正确的结果，但是当您使用 cuda-memcheck 时，会报告错误。您有一个会导致此问题的一般设计问题。

在 C++ 中，当对象在花括号区域内定义时：

{
  {
    Object A;
    // object A is in-scope here
  }
  // object A is out-of-scope here
}
// object A is out of scope here
k<<<...>>>(anything that points to something in object A); // is illegal

然后您退出该区域，该区域内定义的对象现在超出范围。对于具有构造函数/析构函数的对象，当它超出范围时，这通常意味着the destructor of the object will be called。对于thrust::device_vector（或std::vector），这将释放与该向量关联的任何底层存储。这不一定会“擦除”任何数据，但尝试使用该数据是非法的，在 C++ 中将被视为 UB（未定义行为）。

当您在范围内区域内建立指向此类数据的指针，然后超出范围时，这些指针不再指向可以合法访问的任何内容，因此尝试取消引用指针将是非法的/ UB。您的代码正在执行此操作。 是的，它似乎给出了正确的答案，因为在释放时实际上并没有删除任何内容，但是代码设计是非法的，cuda-memcheck 会突出显示这一点。

我想一个解决方法是将所有这些东西从内部花括号中拉出，并将其放在main 范围内，就像d_nFluidNeighborsCrossFluids device_vector 一样。但您可能还想重新考虑您的一般数据组织策略并扁平化您的数据。