摘要
阵[A - B - - - C]在设备内存中，但希望[ABC] -有什么用CUDA C的最快方法？
上下文
我有一个数组A整数对设备（GPU）存储器。 在每次迭代中，我随机选择是大于0的几个元件，并从中减去1。 我维持排序查找数组L那些等于0的元素：
Array A:
@ iteration i: [0 1 0 3 3 2 0 1 2 3]
@ iteration i + 1: [0 0 0 3 2 2 0 1 2 3]
Lookup for 0-elements L:
@ iteration i: [0 - 2 - - - 6 - - -] -> want compacted form: [0 2 6]
@ iteration i + 1: [0 1 2 - - - 6 - - -] -> want compacted form: [0 1 2 6]
（ 在这里，我随机选择元件1和4从减去1。以我在CUDA C实现中，每个线程映射到在元件A ，因此查找数组是稀疏，以防止数据种族和保持排序的排序（例如[0 1 2 6]而不是[0 2 6 1]
后来，我将只对那些对于等于0。因此，我需要用于压缩稀疏查找阵列元件做一些操作L ，这样我可以映射到线程0元素。
因此，什么是压缩上设备存储器中的稀疏数组具有CUDA C的最有效的方法是什么？
非常感谢。
假设我有：
int V[] = {1, 2, 0, 0, 5};
而我期望的结果是：
int R[] = {1, 2, 5}
实际上，我们在删除那个是零，或复制的元素只有当非零元素。
#include <thrust/device_ptr.h>
#include <thrust/copy.h>
#include <stdio.h>
#define SIZE 5
#define cudaCheckErrors(msg) \
do { \
cudaError_t __err = cudaGetLastError(); \
if (__err != cudaSuccess) { \
fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \
msg, cudaGetErrorString(__err), \
__FILE__, __LINE__); \
fprintf(stderr, "*** FAILED - ABORTING\n"); \
exit(1); \
} \
} while (0)
struct is_not_zero
{
__host__ __device__
bool operator()(const int x)
{
return (x != 0);
}
};
int main(){
int V[] = {1, 2, 0, 0, 5};
int R[] = {0, 0, 0, 0, 0};
int *d_V, *d_R;
cudaMalloc((void **)&d_V, SIZE*sizeof(int));
cudaCheckErrors("cudaMalloc1 fail");
cudaMalloc((void **)&d_R, SIZE*sizeof(int));
cudaCheckErrors("cudaMalloc2 fail");
cudaMemcpy(d_V, V, SIZE*sizeof(int), cudaMemcpyHostToDevice);
cudaCheckErrors("cudaMemcpy1 fail");
thrust::device_ptr<int> dp_V(d_V);
thrust::device_ptr<int> dp_R(d_R);
thrust::copy_if(dp_V, dp_V + SIZE, dp_R, is_not_zero());
cudaMemcpy(R, d_R, SIZE*sizeof(int), cudaMemcpyDeviceToHost);
cudaCheckErrors("cudaMemcpy2 fail");
for (int i = 0; i<3; i++)
printf("R[%d]: %d\n", i, R[i]);
return 0;
}
该结构确定指标为我们提供了对于零个元素测试仿函数。 请注意，在推力，没有内核，我们不直接写入设备代码。 所有这一切发生在幕后。 我肯定会建议自己熟悉与快速入门指南 ，以免把这个问题变成上推力的教程。
审查意见后，我觉得代码的这个修改后的版本将解决该CUDA 4.0的问题：
#include <thrust/device_ptr.h>
#include <thrust/copy.h>
#include <thrust/device_vector.h>
#include <thrust/host_vector.h>
#include <stdio.h>
#define SIZE 5
struct is_not_zero
{
__host__ __device__
bool operator()(const int x)
{
return (x != 0);
}
};
int main(){
int V[] = {1, 2, 0, 0, 5};
int R[] = {0, 0, 0, 0, 0};
thrust::host_vector<int> h_V(V, V+SIZE);
thrust::device_vector<int> d_V = h_V;
thrust::device_vector<int> d_R(SIZE, 0);
thrust::copy_if(d_V.begin(), d_V.end(), d_R.begin(), is_not_zero());
thrust::host_vector<int> h_R = d_R;
thrust::copy(h_R.begin(), h_R.end(), R);
for (int i = 0; i<3; i++)
printf("R[%d]: %d\n", i, R[i]);
return 0;
}
