title: Generate Cartesian Product using more than two lists on GPU
body: <p>I would like to know how to generate a Cartesian product of more than two lists using CUDA.</p>
  <p>How do I make this code work with three or more lists?</p>
  <p>It works with two lists but not with three, I tried /, % without success.</p>
  <p>Its basic.</p>
  <pre><code>#include &lt;thrust/device_vector.h&gt;
      #include &lt;thrust/pair.h&gt;
      #include &lt;thrust/copy.h&gt;
      #include &lt;iterator&gt;
      
      __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 &lt; 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(&quot;a[a_idx] and b[b_idx] and c[c_idx] are: %d %d %d\n\n&quot;,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&lt;int&gt; a(2);
        
        // initialize individual elements
        a[0] = 1; 
        a[1] = 2; 
       
      
        // b has storage for 2 integers
        thrust::device_vector&lt;int&gt; b(2);
        
        // initialize individual elements
        b[0] = 3; 
        b[1] = 4; 
       
      
         // d has storage for 2 integers
        thrust::device_vector&lt;int&gt; 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 &quot;raw&quot; pointer from a pointer-like type, simply returning the wrapped pointer, should it exist.
        cartesian_product&lt;&lt;&lt;num_blocks, block_size&gt;&gt;&gt;(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;
      }
  </code></pre>
  <p>How do I get the right c_idx in the kernel and subsequent arrays if I want more than three lists?</p>