title: Optimize CUDA with Thrust in a loop
body: <p>Given the following piece of code, generating a kind of code dictionary with CUDA using thrust (C++ template library for CUDA):</p>
  
  <pre><code>thrust::device_vector&lt;float&gt; dCodes(codes-&gt;begin(), codes-&gt;end());
  thrust::device_vector&lt;int&gt; dCounts(counts-&gt;begin(), counts-&gt;end());
  thrust::device_vector&lt;int&gt; newCounts(counts-&gt;size());
  
  for (int i = 0; i &lt; 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] &gt; 0) {
          newCounts[i]--;
      }
  
      //Remove
      thrust::detail::normal_iterator&lt;thrust::device_ptr&lt;float&gt; &gt; newEnd = thrust::remove(dCodes.begin()+i+1, dCodes.end(), code);
      int dist = thrust::distance(dCodes.begin(), newEnd);
      dCodes.resize(dist);
      newCounts.resize(dist);
  }
  
  codes-&gt;resize(dCodes.size());
  counts-&gt;resize(newCounts.size());
  
  thrust::copy(dCodes.begin(), dCodes.end(), codes-&gt;begin());
  thrust::copy(newCounts.begin(), newCounts.end(), counts-&gt;begin());
  </code></pre>
  
  <p>The problem is, that i've noticed multiple copies of 4 bytes, by using CUDA visual profiler.  IMO this is generated by</p>
  
  <ol>
  <li>The loop counter <em>i</em></li>
  <li><em>float code</em>, <em>int count</em> and <em>dist</em></li>
  <li>Every access to <em>i</em> and the variables noted above</li>
  </ol>
  
  <p>This seems to slow down everything (sequential copying of 4 bytes is no fun...).</p>
  
  <p>So, how i'm telling thrust, that these variables shall be handled on the device? Or are they already?</p>
  
  <p>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).</p>