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

Question

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

我如何使这段代码工作与三个或更多的列表?

它可以处理两个列表，但不能处理三个列表，I tried/， % without success.

其基本的。

#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 %dnn",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)
{
// The first dimension of the grid can be much bigger than the other
// two, so it is generally the best fit for lexicographic indexing.
// On the other hand you might want handle this differently
// for coalescing or when you know size limits beforehand.
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>>>(/* ... */);
/* ... */
}