查找 CUDA 数组中每个不同值的第一个索引
Finding the first index of every distinct value in CUDA array
假设我们有一个这样的数组:
0, 0, 0, 1, 2, 2, 2, 3, 3, 4, ...
我想要每个值第一次出现的索引,所以在这个例子中是 [0, 3, 4, 7, 9]。该数组已排序,所有可能的值都是已知且连续的。
我可能的解决方案是为这个数组中的每个元素使用一个内核,并使用 atomicmin 来保存最低索引。但我认为可能有更好的方法。
正如@tera 指出的那样,您可以将一个数字与前一个数字进行比较,以确定它是否是唯一数字序列中的第一个出现的数字。您可以编写一个内核来为这个条件生成一个掩码,这样掩码数组包含一个数字的索引,该数字是第一次出现,否则是一个负数(比如 -1,因为它不能是索引)。之后,使用推力通过谓词计算非-1 值。然后使用与上面相同的谓词从掩码中复制这些值。最后将结果复制回host。
这是上述方法的示例实现。
#include <iostream>
#include <cuda_runtime.h>
#include <thrust/device_vector.h>
#include <thrust/count.h>
#include <thrust/copy.h>
using namespace std;
//Copy index
__global__ void is_first_occurence(int* input, int* is, int count)
{
const int tid = blockIdx.x * blockDim.x + threadIdx.x;
if(tid<count)
{
if(tid == 0)
{
is[0] = 0;
}
else if(input[tid] != input[tid-1])
{
is[tid] = tid;
}
else
{
is[tid] = -1;
}
}
}
struct isFirst
{
__host__ __device__ bool operator()(const int x)
{
return (x != -1);
}
};
int main(int argc, char** argv)
{
const int count = 13;
std::vector<int> arr = { 0, 0, 0, 1, 1, 2, 2, 2, 3, 3, 4, 4 ,4 };
thrust::device_vector<int> arr_d = arr;
thrust::device_vector<int> mask_d(arr_d.size());
int* pArr = thrust::raw_pointer_cast(arr_d.data() );
int* pMask = thrust::raw_pointer_cast(mask_d.data() );
dim3 block(16);
dim3 grid((count + block.x -1)/block.x);
is_first_occurence<<<grid,block>>>(pArr, pMask, count);
cudaDeviceSynchronize();
int num_unique = thrust::count_if(mask_d.begin(), mask_d.end(), isFirst());
thrust::copy_if(mask_d.begin(), mask_d.end(), arr_d.begin(), isFirst());
std::vector<int> unique_indices(num_unique);
thrust::copy(arr_d.begin(), arr_d.begin() + num_unique, unique_indices.begin());
for(auto i:unique_indices)
{
cout<<i<<endl;
}
return 0;
}
使用以下命令编译和测试:
nvcc -o get_unique get_unique.cu -std=c++11 -arch=sm_61
如果您提供一个索引向量,例如通过 thrust::sequence()
。这是一个有效的例子:
$ cat t3.cu
#include <thrust/device_vector.h>
#include <thrust/copy.h>
#include <thrust/unique.h>
#include <thrust/sequence.h>
#include <iostream>
int main(){
int keys[] = {0, 0, 0, 1, 2, 2, 2, 3, 3, 4};
int ks = sizeof(keys)/sizeof(keys[0]);
thrust::device_vector<int> d_keys(keys, keys+ks);
thrust::device_vector<int> d_result(ks);
thrust::sequence(d_result.begin(), d_result.end());
int rs = (thrust::unique_by_key(d_keys.begin(), d_keys.end(), d_result.begin())).first - d_keys.begin();
thrust::copy_n(d_result.begin(), rs, std::ostream_iterator<int>(std::cout, ","));
std::cout << std::endl;
}
$ nvcc -arch=sm_35 -o t3 t3.cu
$ ./t3
0,3,4,7,9,
$
这里发生的重要 activity 是 流压缩 并且 thrust 为各种用例提供了 nice set of routines。例如,此操作也可以使用 thrust::unique_copy()
完成,在这种情况下,如果代码复杂一些,您可以消除对 thrust::sequence()
调用的需要(它将被 thrust::counting_iterator
替换)与您的数据和适当的选择函子压缩在一起),但它仍然需要相同长度的输出向量。
假设我们有一个这样的数组:
0, 0, 0, 1, 2, 2, 2, 3, 3, 4, ...
我想要每个值第一次出现的索引,所以在这个例子中是 [0, 3, 4, 7, 9]。该数组已排序,所有可能的值都是已知且连续的。
我可能的解决方案是为这个数组中的每个元素使用一个内核,并使用 atomicmin 来保存最低索引。但我认为可能有更好的方法。
正如@tera 指出的那样,您可以将一个数字与前一个数字进行比较,以确定它是否是唯一数字序列中的第一个出现的数字。您可以编写一个内核来为这个条件生成一个掩码,这样掩码数组包含一个数字的索引,该数字是第一次出现,否则是一个负数(比如 -1,因为它不能是索引)。之后,使用推力通过谓词计算非-1 值。然后使用与上面相同的谓词从掩码中复制这些值。最后将结果复制回host。
这是上述方法的示例实现。
#include <iostream>
#include <cuda_runtime.h>
#include <thrust/device_vector.h>
#include <thrust/count.h>
#include <thrust/copy.h>
using namespace std;
//Copy index
__global__ void is_first_occurence(int* input, int* is, int count)
{
const int tid = blockIdx.x * blockDim.x + threadIdx.x;
if(tid<count)
{
if(tid == 0)
{
is[0] = 0;
}
else if(input[tid] != input[tid-1])
{
is[tid] = tid;
}
else
{
is[tid] = -1;
}
}
}
struct isFirst
{
__host__ __device__ bool operator()(const int x)
{
return (x != -1);
}
};
int main(int argc, char** argv)
{
const int count = 13;
std::vector<int> arr = { 0, 0, 0, 1, 1, 2, 2, 2, 3, 3, 4, 4 ,4 };
thrust::device_vector<int> arr_d = arr;
thrust::device_vector<int> mask_d(arr_d.size());
int* pArr = thrust::raw_pointer_cast(arr_d.data() );
int* pMask = thrust::raw_pointer_cast(mask_d.data() );
dim3 block(16);
dim3 grid((count + block.x -1)/block.x);
is_first_occurence<<<grid,block>>>(pArr, pMask, count);
cudaDeviceSynchronize();
int num_unique = thrust::count_if(mask_d.begin(), mask_d.end(), isFirst());
thrust::copy_if(mask_d.begin(), mask_d.end(), arr_d.begin(), isFirst());
std::vector<int> unique_indices(num_unique);
thrust::copy(arr_d.begin(), arr_d.begin() + num_unique, unique_indices.begin());
for(auto i:unique_indices)
{
cout<<i<<endl;
}
return 0;
}
使用以下命令编译和测试:
nvcc -o get_unique get_unique.cu -std=c++11 -arch=sm_61
如果您提供一个索引向量,例如通过 thrust::sequence()
。这是一个有效的例子:
$ cat t3.cu
#include <thrust/device_vector.h>
#include <thrust/copy.h>
#include <thrust/unique.h>
#include <thrust/sequence.h>
#include <iostream>
int main(){
int keys[] = {0, 0, 0, 1, 2, 2, 2, 3, 3, 4};
int ks = sizeof(keys)/sizeof(keys[0]);
thrust::device_vector<int> d_keys(keys, keys+ks);
thrust::device_vector<int> d_result(ks);
thrust::sequence(d_result.begin(), d_result.end());
int rs = (thrust::unique_by_key(d_keys.begin(), d_keys.end(), d_result.begin())).first - d_keys.begin();
thrust::copy_n(d_result.begin(), rs, std::ostream_iterator<int>(std::cout, ","));
std::cout << std::endl;
}
$ nvcc -arch=sm_35 -o t3 t3.cu
$ ./t3
0,3,4,7,9,
$
这里发生的重要 activity 是 流压缩 并且 thrust 为各种用例提供了 nice set of routines。例如,此操作也可以使用 thrust::unique_copy()
完成,在这种情况下,如果代码复杂一些,您可以消除对 thrust::sequence()
调用的需要(它将被 thrust::counting_iterator
替换)与您的数据和适当的选择函子压缩在一起),但它仍然需要相同长度的输出向量。