Databases
StackOverflow2013
Postsdbo
POIncreasing achieved occupancy doesn't enhance computation speed linearly
 

Note that there are some explanatory texts on larger screens.

plurals
  1. POIncreasing achieved occupancy doesn't enhance computation speed linearly
    primarykey
    data
    text
    <p>I had a CUDA program in which kernel registers were limiting <strong>maximum theoretical achieved occupancy</strong> to %50. So I decided to use shared memory instead of registers for those variables that were constant between block threads and were almost read-only throughout kernel run. I cannot provide source code here; what I did was conceptually like this:</p> <p>My initial program:</p> <pre><code>__global__ void GPU_Kernel (...) { __shared__ int sharedData[N]; //N:maximum amount that doesn't limit maximum occupancy int r_1 = A; //except for this first initialization, these registers don't change anymore int r_2 = B; ... int r_m = Y; ... //rest of kernel; } </code></pre> <p>I changed above program to:</p> <pre><code>__global__ void GPU_Kernel (...) { __shared__ int sharedData[N-m]; __shared__ int r_1, r_2, ..., r_m; if ( threadIdx.x == 0 ) { r_1 = A; r_2 = B; ... r_m = Y; //last of them } __syncthreads(); ... //rest of kernel } </code></pre> <p>Now threads of warps inside a block perform broadcast reads to access newly created shared memory variables. At the same time, threads don't use too much registers to limit achieved occupancy.</p> <p>The second program has <strong>maximum theoretical achieved occupancy</strong> equal to %100. In actual runs, the average achieved occupancy for the first programs was ~%48 and for the second one is around ~%80. But the issue is enhancement in net speed up is around %5 to %10, much less than what I was anticipating considering improved gained occupancy. Why isn't this correlation linear?</p> <p>Considering below image from Nvidia whitepaper, what I've been thinking was that when achieved occupancy is %50, for example, half of SMX (in newer architectures) cores are idle at a time because excessive requested resources by other cores stop them from being active. Is my understanding flawed? Or is it incomplete to explain above phenomenon? Or is it added <code>__syncthreads();</code> and shared memory accesses cost?</p> <p><img src="https://i.stack.imgur.com/9GF9j.png" alt="enter image description here"></p>
    singulars
    1. This table or related slice is empty.
    1. This table or related slice is empty.
    1. PTQuestion
    1. USFarzad
    1. USFarzad
    plurals
    1. This table or related slice is empty.
    1. This table or related slice is empty.
    1. This table or related slice is empty.
    1. PO
      singulars
      1. PTAnswer
    1. This table or related slice is empty.
    1. This table or related slice is empty.
 

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