CODEWITHSUNDEEP
optimizationcudagpugpgpu
RSFalcon7
RSFalcon7

Reputation: 2311

How to efficiently gather data from threads in CUDA?

I have a application that solves a system of equations in CUDA, I know for sure that each thread can find up to 4 solutions, but how can I copy then back to the host?

I'm passing a huge array with enough space to all threads store 4 solutions (4 doubles for each solution), and another one with the number of solutions per thread, however that's a naive solution, and is the current bottleneck of my kernel.

I really like to optimize this. The main problem is concatenate a variable number of solutions per thread in a single array.

Upvotes: 5

Views: 1483

Answers (1)

Roger Dahl
Roger Dahl

Reputation: 15734

The functionality you're looking for is called stream compaction.

You probably do need to provide an array that contains room for 4 solutions per thread because attempting to directly store the results in a compact form is likely to create so many dependencies between the threads that the performance gained in being able to copy less data back to the host is lost by a longer kernel execution time. The exception to this is if almost all of the threads find no solutions. In that case, you might be able to use an atomic operation to maintain an index into an array. So, for each solution that is found, you would store it in an array at an index and then use an atomic operation to increase the index. I think it would be safe to use atomicAdd() for this. Before storing a result, the thread would use atomicAdd() to increase the index by one. atomicAdd() returns the old value, and the thread can store the result using the old value as the index.

However, given a more common situation, where there's a fair number of results, the best solution will be to perform a compacting operation as a separate step. One way to do this is with thrust::copy_if. See this question for some more background.

Upvotes: 4

Related Questions