在 GPU 上使用两个以上列表生成笛卡尔积

Question

我想知道如何使用 CUDA 生成两个以上列表的笛卡尔积。

如何使这段代码适用于三个或更多列表？

它适用于两个列表，但不适用于三个列表，我尝试了 /, % 没有成功。

它是基本的。

#include <thrust/device_vector.h>
    #include <thrust/pair.h>
    #include <thrust/copy.h>
    #include <iterator>
    
    __global__ void cartesian_product(const int *a, size_t a_size,
                                      const int *b, size_t b_size,
                                      const int *c, size_t c_size)
    {
      unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;
    
      if(idx < a_size * b_size * c_size) 
      {
        unsigned int a_idx = idx / a_size;
        unsigned int b_idx = idx % a_size;
        
        // ? 
        unsigned int c_idx = idx % a_size;
    
    
        
        printf("a[a_idx] and b[b_idx] and c[c_idx] are: %d %d %d\n\n",a[a_idx], b[b_idx], c[c_idx]);
        //1 3 5 , 1 3 6 , 1 4 5 , 1 4 6 , 2 3 5 , 2 3 6 , 2 4 5 , 2 4 6  
        //0 0 0 , 0 0 1 , 0 1 0 , 0 1 1 , 1 0 0 , 1 0 1 , 1 1 0 , 1 1 1
      }
    }
    
    int main()
    {
      
      
      // host_vector is stored in host memory while device_vector livesin GPU device memory.
      // a has storage for 2 integers
      thrust::device_vector<int> a(2);
      
      // initialize individual elements
      a[0] = 1; 
      a[1] = 2; 
     
    
      // b has storage for 2 integers
      thrust::device_vector<int> b(2);
      
      // initialize individual elements
      b[0] = 3; 
      b[1] = 4; 
     
    
       // d has storage for 2 integers
      thrust::device_vector<int> c(2);
      
      // initialize individual elements
      c[0] = 5; 
      c[1] = 6;
      
       
      unsigned int block_size = 256;
      unsigned int num_blocks = (8 + (block_size - 1)) / block_size;
    
    
      // raw_pointer_cast creates a "raw" pointer from a pointer-like type, simply returning the wrapped pointer, should it exist.
      cartesian_product<<<num_blocks, block_size>>>(thrust::raw_pointer_cast(a.data()), a.size(),
                                                    thrust::raw_pointer_cast(b.data()), b.size(),
                                                    thrust::raw_pointer_cast(c.data()), c.size());

      
      
      return 0;
    }

如果我想要三个以上的列表，如何在内核和后续数组中获取正确的 c_idx？

Answer 1

在我看来，您想要“词法索引”：

idx == (a_idx * b_size + b_idx) * c_size + c_idx

所以你得到这样的索引：

c_idx = idx % c_size;
b_idx = (idx / c_size) % b_size;
a_idx = (idx / c_size) / b_size;

这很容易推广到更多维度。例如。你有四个维度

idx == ((a_idx * b_size + b_idx) * c_size + c_idx) * d_size + d_idx

然后：

d_idx = idx % d_size;
c_idx = (idx / d_size) % c_size;
b_idx = ((idx / d_size) / c_size) % b_size;
a_idx = ((idx / d_size) / c_size) / b_size;

在 C/C++ 编程中，人们喜欢使用它来计算索引到表示多维数据集的一维动态数组中。在 CUDA 中，您通常不需要那么多，因为 CUDA 最多为您提供三维threadIdx/blockIdx/等。所以对于三个数组的笛卡尔积，您不需要这种技术，但是可以只使用内在的 CUDA 功能。即使在三个以上的情况下，性能最高的解决方案也会从内核的三个维度中的两个维度中获取两个索引，并在第三个维度上使用词法索引：

__global__ void cartesian_product_5d(const int *a, size_t a_size,
                                  const int *b, size_t b_size,
                                  const int *c, size_t c_size,
                                  const int *d, size_t d_size,
                                  const int *e, size_t e_size)
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    int d_idx = blockIdx.y * blockDim.y + threadIdx.y;
    int e_idx = blockIdx.z * blockDim.z + threadIdx.z;
    /* idx == (c_idx * b_size + b_idx) * a_size + a_idx */
    int a_idx = idx % a_size;
    int b_idx = (idx / a_size) % b_size;
    int c_idx = (idx / a_size) / b_size;

    /* ... */
}
 
int main()
{
    /* ... */
    dim3 threadsPerBlock(8, 8, 8);
    dim3 numBlocks((a_size + b_size + c_size + threadsPerBlock.x - 1) /
                   threadsPerBlock.x,
                   (d_size + threadsPerBlock.y - 1) / threadsPerBlock.y,
                   (e_size + threadsPerBlock.z - 1) / threadsPerBlock.z);
    cartesian_product_5d<<<numBlocks, threadsPerBlock>>>(/* ... */);
    /* ... */
}