Optimize CUDA with Thrust in a loop
Posted
by macs
on Stack Overflow
See other posts from Stack Overflow
or by macs
Published on 2010-03-08T22:16:58Z
Indexed on
2010/03/08
22:21 UTC
Read the original article
Hit count: 349
Given the following piece of code, generating a kind of code dictionary with CUDA using thrust (C++ template library for CUDA):
thrust::device_vector<float> dCodes(codes->begin(), codes->end());
thrust::device_vector<int> dCounts(counts->begin(), counts->end());
thrust::device_vector<int> newCounts(counts->size());
for (int i = 0; i < dCodes.size(); i++) {
float code = dCodes[i];
int count = thrust::count(dCodes.begin(), dCodes.end(), code);
newCounts[i] = dCounts[i] + count;
//Had we already a count in one of the last runs?
if (dCounts[i] > 0) {
newCounts[i]--;
}
//Remove
thrust::detail::normal_iterator<thrust::device_ptr<float> > newEnd = thrust::remove(dCodes.begin()+i+1, dCodes.end(), code);
int dist = thrust::distance(dCodes.begin(), newEnd);
dCodes.resize(dist);
newCounts.resize(dist);
}
codes->resize(dCodes.size());
counts->resize(newCounts.size());
thrust::copy(dCodes.begin(), dCodes.end(), codes->begin());
thrust::copy(newCounts.begin(), newCounts.end(), counts->begin());
The problem is, that i've noticed multiple copies of 4 bytes, by using CUDA visual profiler. IMO this is generated by
- The loop counter i
- float code, int count and dist
- Every access to i and the variables noted above
This seems to slow down everything (sequential copying of 4 bytes is no fun...).
So, how i'm telling thrust, that these variables shall be handled on the device? Or are they already?
Using thrust::device_ptr seems not sufficient for me, because i'm not sure whether the for loop around runs on host or on device (which could also be another reason for the slowliness).
© Stack Overflow or respective owner