Cuda 内核返回向量
Posted
技术标签:
【中文标题】Cuda 内核返回向量【英文标题】:Cuda kernel returning vectors 【发布时间】:2014-03-14 05:38:52 【问题描述】:我有一个单词列表,我的目标是匹配一个非常长的短语中的每个单词。 我在匹配每个单词时没有问题,我唯一的问题是返回一个包含每个匹配信息的结构向量。
在代码中:
typedef struct
int A, B, C; Match;
__global__ void Find(veryLongPhrase * _phrase, Words * _word_list, vector<Match> * _matches)
int a, b, c;
[...] //Parallel search for each word in the phrase
if(match) //When an occurrence is found
_matches.push_back(new Match A = a, B = b, C = c ); //Here comes the unknown, what should I do here???
main()
[...]
veryLongPhrase * myPhrase = "The quick brown fox jumps over the lazy dog etc etc etc..."
Words * wordList = "the", "lazy";
vector<Match> * matches; //Obviously I can't pass a vector to a kernel
Find<<< X, Y >>>(myPhrase, wordList, matches);
[...]
我已经尝试过 Thrust 库,但没有任何成功,你能建议我任何解决方案吗?
非常感谢。
【问题讨论】:
推力分区似乎很适合解决这个问题。 虽然我应该指出,这将是一个关于它是否比 CPU 快的废话。内存布局将是 AWFUL(我假设您将其布置为 GPU 内存中的 char**)。分区并不是您在美好的一天可以在 GPU 上执行的最快操作。打开和关闭复制数据将是一场噩梦(对于内存复制循环)。如果您的短语不超过几千字,我认为它不会加快您的代码速度 【参考方案1】:这样的东西应该可以工作(在浏览器中编码,未经测试):
// N is the maximum number of structs to insert
#define N 10000
typedef struct
int A, B, C; Match;
__device__ Match dev_data[N];
__device__ int dev_count = 0;
__device__ int my_push_back(Match * mt)
int insert_pt = atomicAdd(&dev_count, 1);
if (insert_pt < N)
dev_data[insert_pt] = *mt;
return insert_pt;
else return -1;
__global__ void Find(veryLongPhrase * _phrase, Words * _word_list, vector<Match> * _matches)
int a, b, c;
[...] //Parallel search for each word in the phrase
if(match) //When an occurrence is found
my_push_back(new Match A = a, B = b, C = c );
main()
[...]
veryLongPhrase * myPhrase = "The quick brown fox jumps over the lazy dog etc etc etc..."
Words * wordList = "the", "lazy";
Find<<< X, Y >>>(myPhrase, wordList);
int dsize;
cudaMemcpyFromSymbol(&dsize, dev_count, sizeof(int));
vector<Match> results(dsize);
cudaMemcpyFromSymbol(&(results[0]), dev_data, dsize*sizeof(Match));
[...]
这需要原子操作的计算能力 1.1 或更高。
nvcc -arch=sm_11 ...
这是一个有效的例子:
$ cat t347.cu
#include <iostream>
#include <vector>
// N is the maximum number of structs to insert
#define N 10000
typedef struct
int A, B, C; Match;
__device__ Match dev_data[N];
__device__ int dev_count = 0;
__device__ int my_push_back(Match & mt)
int insert_pt = atomicAdd(&dev_count, 1);
if (insert_pt < N)
dev_data[insert_pt] = mt;
return insert_pt;
else return -1;
__global__ void Find()
if(threadIdx.x < 10) //Simulate a found occurrence
Match a = .A = 1, .B = 2, .C = 3 ;
my_push_back(a);
main()
Find<<< 2, 256 >>>();
int dsize;
cudaMemcpyFromSymbol(&dsize, dev_count, sizeof(int));
if (dsize >= N) printf("overflow error\n"); return 1;
std::vector<Match> results(dsize);
cudaMemcpyFromSymbol(&(results[0]), dev_data, dsize*sizeof(Match));
std::cout << "number of matches = " << dsize << std::endl;
std::cout << "A = " << results[dsize-1].A << std:: endl;
std::cout << "B = " << results[dsize-1].B << std:: endl;
std::cout << "C = " << results[dsize-1].C << std:: endl;
$ nvcc -arch=sm_11 -o t347 t347.cu
$ ./t347
number of matches = 20
A = 1
B = 2
C = 3
$
请注意,在这种情况下,我的 Match
结果结构创建是不同的,我是通过引用传递的,但概念是相同的。
【讨论】:
我的情况非常相似,这个解决方案效果很好。感谢罗伯特的回答。 只是想澄清两件事:1)MS VS 编辑器总是会抱怨atomicAdd
,因为它直到编译时才定义。 2) 同样,它会抱怨cudaMemcpyFromSymbol
的第二个参数,因为函数模板显示const void *
,但不支持指向设备内存的指针。试图摆脱警告无济于事。以上是关于Cuda 内核返回向量的主要内容,如果未能解决你的问题,请参考以下文章