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

Question

我有一个问题，我在其中找到了很多线索，但 none 明确回答了我的问题。我正在尝试使用推力在 GPU 内核中创建一个多维数组。展平会很困难，因为所有维度都是非均匀的，我会升至 4D。现在我知道我不能有 device_vectors of 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 会工作，但是 return 错误的值。我这样访问它（同样，从内核中）：

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

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

谢谢你抽出时间，祝你今天愉快。

编辑：这是一个最小的可重现示例。由于某种原因，尽管与我的代码具有相同的结构，但这些值似乎是正确的，但 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；你的既不是最小的，也不是完整的，也不是可验证的。

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

I know I cannot have device_vectors of device_vectors, for whichever underlying reason (explanation would be welcome)

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

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

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

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

Answer 2

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

d_nFluidNeighborsCrossFluids[...]...;

是方括号运算符。该运营商是主机运营商（仅）。它在设备代码中不可用。像这样的问题导致一般性的说法是 "thrust::device_vector is not usable in device code." device_vector object 本身一般是不可用的。但是，如果您尝试通过原始指针访问它，它包含的数据可用于设备代码。

这是一个推力设备向量的示例，其中包含指向其他设备向量中包含的数据的指针数组。只要您不尝试使用 thrust::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

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

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

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

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

Why does thrust::device_vector not seem to have a chance to hold raw pointers to other device_vectors?

cuda

multidimensional-array

thrust