时间戳

首页 > 解决方案 > 为什么推力::device_vector 似乎没有机会保存指向其他 device_vectors 的原始指针？

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

问题描述

我有一个问题，我发现了很多线程，但没有一个明确回答我的问题。我正在尝试使用推力在 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();
}

标签： multidimensional-arraycudathrust

解决方案

你真的应该提供一个最小的、完整的、可验证/可重现的例子；你的既不是最小的，也不是完整的，也不是可验证的。

但是，我将回答您的附带问题：

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

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

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

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

最后，请记住，指针追逐可能会成为性能杀手，尤其是在 GPU 上。您可能需要考虑在某种程度上按摩您的数据结构，以使其易于扁平化。

推荐阅读