根据包含全局地址的各自索引向量将两个不同长度的向量拼接到具有推力的公共长度的新向量上

标记

这个问题已经困扰了我好几年了。我从这个论坛中学到了很多c ++和cuda。以前,我在fortran串行代码中使用许多条件语句编写了以下内容,并使用gotos,因为我找不到执行此操作的聪明方法。
这是问题所在。

给定4个向量:

int indx(nshape);
float dnx(nshape); 
/* nshape > nord */
int indy(nord);
float dny(nord);

indx和indy是包含全局坐标的索引向量(分别为值dnx,dny的键)。在解析此所需的隔行扫描/拼接功能之前,尚不知道它们的全局范围。所有已知的是,可能的局部范围的长度可以为[0,nord * nord],向量indx和indy中的最大值和最小值。

我想创建长度相同的新向量dn1和dn2,其中包含dnx和dny的原始值,但是将其扩展为用零填充原始向量dnx和dny,以便它们不包含其他向量的所有全局坐标。它们将形成需要全局地址对齐的外部产品的向量。

我无法在网上找到任何有关在c ++中使用逻辑掩码(例如在fortran中进行并行化)的参考。我的出发点是使用推力库stable_sort升序,使用binary_search比较数组,分区等。也许有一种清晰简洁的方法。

下面的示例索引和值vecs通常不会从0开始或与临时索引向量的本地寻址一致,也不会与任何偶数奇数模式重合-这些值仅用于帮助说明。)

indx[]={0,2,4,6,8,10,12};  indy[]={1, 2, 3, 4};
dnx[]={99,99,99,99,99,99,99};  dny[]={66,66,66,66};

ind[]={0,1,2,3,4,6,8,10,12}
dn1[]={99,0,99,0,99,99,99,99,99}
dn2[]={0,66,66,66,66,0,0,0,0}

以前我做过类似下面的事情,其中​​内核根据以下条件应用比较,填充和流,然后继续通过这些条件行之一再次返回,直到最大局部索引超过最大向量i,ei,j的长度。 > nshape:

3
if(indx[i] < indy[j]{kernel_1; i++; if(i > nshape){return}; goto 3}
if(indx[i] == indy[j]){kernel_2;i++;j++; if(i || j > nshape) {return}; goto 3}
if(indx[i] > indy[j]{kernel_3, j++, if(j>nshape){return}; goto 3}

对不起,杂种伪代码。我真的很期待使用c ++,cuda,推力的任何想法或更好的解决方案。非常感谢。标记

罗伯特·克罗维拉

我想到的一种方法是进行并行二进制搜索,从全局索引向量中获取每个值,然后查看键向量中是否有匹配项。如果在键向量中有匹配项,则将相应的值放在结果中。如果键向量中没有匹配项,则将0放入结果中。

因此,对于以下位置中的每个位置:

ind[]={0,1,2,3,4,6,8,10,12}

查看以下位置是否有匹配的索引:

indy[]={1, 2, 3, 4};

我们将使用并行二进制搜索来执行此操作,并返回(中的匹配值的indy对应索引

如果确实找到匹配项,则在结果中所讨论的位置处,我们将放置对应于的值indy[matching_index],即。dny[matching_index]否则,将结果置零。

就推力而言,我可以将其减少为两个推力。

第一个是thrust::lower_bound运算,实际上是向量化/并行二进制搜索。就像在CUDA情况下一样,我们使用二分搜索来获取全局向量(ind)的每个元素,并查看关键向量(例如indx)中是否存在匹配项,并返回关键向量中匹配位置的索引(较低边界)。

第二个调用是的复杂用法thrust::for_each我们extend_functor为该for_each操作创建一个特殊的函子(),该函子使用指向键向量(例如indx的开头,其长度以及指向值向量(例如dnx的指针进行初始化所述extend_functor然后取全局矢量的3元组,下侧结合的矢量,并将结果矢量值,并执行剩余的步骤。如果下界值在关键向量的长度之内,则检查下界是否在关键向量和全局向量之间产生匹配。如果是这样,请将相应的值放入结果向量,否则将零放入结果向量。

后续代码使用CUDA以及推力来实现此目的。

#include <iostream>
#include <thrust/device_vector.h>
#include <thrust/binary_search.h>
#include <thrust/copy.h>

#define MAX_DSIZE 32768
#define nTPB 256


struct extend_functor
{
  int vec_len;
  int *vec1;
  float *vec2;
  extend_functor(int *_vec1, int _vec_len, float *_vec2) : vec1(_vec1), vec2(_vec2),  vec_len(_vec_len) {};
  template <typename Tuple>
  __device__ __host__
  void operator()(const Tuple &my_tuple) {
    float result = 0.0f;
    if (thrust::get<1>(my_tuple) < vec_len)
      if (thrust::get<0>(my_tuple) == vec1[thrust::get<1>(my_tuple)]) result = vec2[thrust::get<1>(my_tuple)];
    thrust::get<2>(my_tuple) = result;
  }

};

// binary search, looking for key in a

__device__ void bsearch_range(const int *a, const int key, const unsigned len_a, unsigned *idx){
  unsigned lower = 0;
  unsigned upper = len_a;
  unsigned midpt;
  while (lower < upper){
    midpt = (lower + upper)>>1;
    if (a[midpt] < key) lower = midpt +1;
    else upper = midpt;
    }
  *idx = lower;
  return;
  }


// k is the key vector
// g is the global index vector
// v is the value vector
// r is the result vector

__global__ void extend_kernel(const int *k, const unsigned len_k, const int *g,  const unsigned len_g, const float *v, float *r){
  unsigned idx = (blockDim.x * blockIdx.x) + threadIdx.x;
  if (idx < len_g) {
    unsigned my_idx;
    int g_key = g[idx];
    bsearch_range(k, g_key, len_k, &my_idx);
    int my_key = -1;
    if (my_idx < len_k)
      my_key = k[my_idx];
    float my_val;
    if (g_key == my_key) my_val = v[my_idx];
    else my_val = 0.0f;
    r[idx] = my_val;
    }
}


int main(){

  int len_x = 7;
  int len_y = 4;
  int len_g = 9;
  int indx[]={0,2,4,6,8,10,12};
  int indy[]={1, 2, 3, 4};
  float dnx[]={91.0f,92.0f,93.0f,94.0f,95.0f,96.0f,97.0f};
  float dny[]={61.0f,62.0f,63.0f,64.0f};
  int ind[]={0,1,2,3,4,6,8,10,12};

  int *h_k, *d_k, *h_g, *d_g;
  float *h_v, *d_v, *h_r, *d_r;

  h_k = (int *)malloc(MAX_DSIZE*sizeof(int));
  h_g = (int *)malloc(MAX_DSIZE*sizeof(int));
  h_v = (float *)malloc(MAX_DSIZE*sizeof(float));
  h_r = (float *)malloc(MAX_DSIZE*sizeof(float));

  cudaMalloc(&d_k, MAX_DSIZE*sizeof(int));
  cudaMalloc(&d_g, MAX_DSIZE*sizeof(int));
  cudaMalloc(&d_v, MAX_DSIZE*sizeof(float));
  cudaMalloc(&d_r, MAX_DSIZE*sizeof(float));

// test case x

  memcpy(h_k, indx, len_x*sizeof(int));
  memcpy(h_g, ind, len_g*sizeof(int));
  memcpy(h_v, dnx, len_x*sizeof(float));

  cudaMemcpy(d_k, h_k, len_x*sizeof(int), cudaMemcpyHostToDevice);
  cudaMemcpy(d_g, h_g, len_g*sizeof(int), cudaMemcpyHostToDevice);
  cudaMemcpy(d_v, h_v, len_x*sizeof(float), cudaMemcpyHostToDevice);
  extend_kernel<<<(len_g+nTPB-1)/nTPB, nTPB >>>(d_k, len_x, d_g, len_g, d_v, d_r);
  cudaMemcpy(h_r, d_r, len_g*sizeof(float), cudaMemcpyDeviceToHost);

  std::cout << "case x result: ";
  for (int i=0; i < len_g; i++)
    std::cout << h_r[i] << " ";
  std::cout << std::endl;


// test case y

  memcpy(h_k, indy, len_y*sizeof(int));
  memcpy(h_g, ind, len_g*sizeof(int));
  memcpy(h_v, dny, len_y*sizeof(float));

  cudaMemcpy(d_k, h_k, len_y*sizeof(int), cudaMemcpyHostToDevice);
  cudaMemcpy(d_g, h_g, len_g*sizeof(int), cudaMemcpyHostToDevice);
  cudaMemcpy(d_v, h_v, len_y*sizeof(float), cudaMemcpyHostToDevice);
  extend_kernel<<<(len_g+nTPB-1)/nTPB, nTPB >>>(d_k, len_y, d_g, len_g, d_v, d_r);
  cudaMemcpy(h_r, d_r, len_g*sizeof(float), cudaMemcpyDeviceToHost);

  std::cout << "case y result: ";
  for (int i=0; i < len_g; i++)
    std::cout << h_r[i] << " ";
  std::cout << std::endl;

// using thrust

  thrust::device_vector<int> tind(ind, ind+len_g);
  thrust::device_vector<int> tindx(indx, indx+len_x);
  thrust::device_vector<float> tdnx(dnx, dnx+len_x);
  thrust::device_vector<float> tresult(len_g);

  thrust::device_vector<int> tbound(len_g);
  thrust::lower_bound(tindx.begin(), tindx.end(), tind.begin(), tind.end(), tbound.begin());
  thrust::for_each(thrust::make_zip_iterator(thrust::make_tuple(tind.begin(), tbound.begin(), tresult.begin())), thrust::make_zip_iterator(thrust::make_tuple(tind.end(), tbound.end(), tresult.end())), extend_functor(thrust::raw_pointer_cast(tindx.data()), len_x, thrust::raw_pointer_cast(tdnx.data())));

  std::cout << "thrust case x result: ";
  thrust::copy(tresult.begin(), tresult.end(), std::ostream_iterator<float>(std::cout, " "));
  std::cout << std::endl;
  return 0;
}

本文收集自互联网,转载请注明来源。

如有侵权,请联系 [email protected] 删除。

编辑于
0

我来说两句

0 条评论
登录 后参与评论

相关文章