Databases
StackOverflow2013
Postsdbo
PO
 

Note that there are some explanatory texts on larger screens.

plurals
  1. PO
    primarykey
    data
    text
    <p>As <a href="https://stackoverflow.com/users/681865/talonmies">talonmies</a> pointed out, you're accessing and processing your data byte-wise, which is far from optimal. A collection of techniques you may want to consider, such as Instruction-Level Parallelism and buffered read/writes, are summarized in the nVidia Webinar <a href="http://nvidia.fullviewmedia.com/gtc2010/0922-a5-2238.html" rel="nofollow noreferrer">Better Performance at Lower Occupancy</a> by Vasily Volkov.</p> <p>In a nutshell, what you want to do is, in each thread, read <em>several</em> <code>uint4</code> in a coalesced way, process them, and only then store them.</p> <p><strong>Update</strong></p> <p>Does it make any difference if you re-write your code as follows?</p> <pre><code>__global__ void logical_and(unsigned int* in, unsigned int* out, int N) { int idx = blockIdx.x*blockDim.x*chunksize+threadIdx.x; unsigned int buff[chunksize]; #pragma unroll for ( int k = 0 ; k &lt; chunksize ; k++ ) buff[k] = in[ blockDim.x*k + idx ]; #pragma unroll for ( int k = 0 ; k &lt; chunksize ; k++ ) buff[k] &amp;= in[ blockDim.x*k + idx + N ]; #pragma unroll for ( int k = 0 ; k &lt; chunksize ; k++ ) out[ blockDim.x*k + idx ] = buff[k]; } </code></pre> <p>Note that I've assumed <code>chunksize</code> is a variable you've <code>#define</code>d somewhere, e.g.</p> <pre><code>#define chunksize 4 </code></pre> <p>And that you have to divide the number of blocks you launch and <code>N</code> by that number. I've also used <code>unsigned int</code> which is just four packed <code>uchar</code>. In your calling function, you may have to cast your pointers accordingly.</p>
    singulars
    1. This table or related slice is empty.
    1. POCUDA kernel is slower than CPU
      singulars
      1. PTQuestion
    1. PTAnswer
    1. USCommunity
    1. USPedro
    plurals
    1. This table or related slice is empty.
    1. This table or related slice is empty.
    1. POCUDA kernel is slower than CPU
      singulars
      1. PTQuestion
    1. This table or related slice is empty.
    1. VO
      singulars
      1. PO
      1. This table or related slice is empty.
      1. VTAcceptedByOriginator
    2. VO
      singulars
      1. PO
      1. This table or related slice is empty.
      1. VTUpMod
    3. VO
      singulars
      1. PO
      1. This table or related slice is empty.
      1. VTUpMod
    1. COgreat webinar, thanks
      singulars
      1. PO
      1. USeel
    2. COI tried process 2 uchar4 in each thread, but the speed remains the same
      singulars
      1. PO
      1. USeel
    3. COUsing uint4 instead of uchar4 was ~15% slower
      singulars
      1. PO
      1. USeel
 

Querying!

 
Guidance
A row detail

Detail views are divided into sections. All the information in the data section comes from columns in the selected row. The other sections display data from other, related rows.

Related data can be related in a to-one or a to-many fashion. Captions of data related in a to-many fashion link to a list view showing a filtered view of the table.

Try moving around until you find a non-empty to-many entry and click on the label to get to one. You can move back to the root by clicking on the database name in the header.

SQuiL has stopped working due to an internal error.

If you are curious you may find further information in the browser console, which is accessible through the devtools (F12).

Reload