Databases
StackOverflow2013
Postsdbo
POCuda reuse of shared memory variable name
 

Note that there are some explanatory texts on larger screens.

plurals
  1. POCuda reuse of shared memory variable name
    primarykey
    data
    text
    <p>I have two cuda kernels that are run, one after the other:</p> <pre><code>__global__ void calculate_histo(const float* const d_logLuminance, unsigned int* d_histogram, float min_logLum, float lumRange, int numBins, int num_elements){ extern __shared__ float sdata[]; int tid = threadIdx.x; int bid = blockIdx.x; int gid = tid * blockDim.x + bid; // load input into __shared__ memory if(gid &lt; num_elements) { sdata[tid] = d_logLuminance[gid]; __syncthreads(); //compute bin value of input int bin = static_cast &lt;int&gt; (floor((d_logLuminance[gid]-min_logLum)/ lumRange * numBins)); //increment histogram at bin value atomicAdd(&amp;(d_histogram[bin]), 1); } } __global__ void blelloch_scan(unsigned int* const d_cdf, unsigned int* d_histogram, int numBins) { extern __shared__ unsigned int sdata[];// allocated on invocation int thid = threadIdx.x; //printf("%i \n", thid); //printf("%i \n", d_histogram[thid]); int offset = 1; sdata[2*thid] = d_histogram[2*thid]; // load input into shared memory sdata[2*thid+1] = d_histogram[2*thid+1]; // build sum in place up the tree for (int d = numBins&gt;&gt;1; d &gt; 0; d &gt;&gt;= 1) { __syncthreads(); if (thid &lt; d) { int ai = offset*(2*thid+1)-1; int bi = offset*(2*thid+2)-1; sdata[bi] += sdata[ai]; } offset *= 2; } if (thid == 0) { sdata[numBins - 1] = 0; } // clear the last element // traverse down tree &amp; build scan for (int d = 1; d &lt; numBins; d *= 2) { offset &gt;&gt;= 1; __syncthreads(); if (thid &lt; d) { int ai = offset*(2*thid+1)-1; int bi = offset*(2*thid+2)-1; float t = sdata[ai]; sdata[ai] = sdata[bi]; sdata[bi] += t; } __syncthreads(); d_cdf[2*thid] = sdata[2*thid]; // write results to device memory d_cdf[2*thid+1] = sdata[2*thid+1]; } } </code></pre> <p>They both use shared memory. The second has an unsigned int array as the shared memory. The first has a float array. I thought I should be able to reuse the same variable name, sdata, for both arrays, since shared memory is cleared after each kernel launch, but I'm getting the error: </p> <pre><code>declaration is incompatible with previous 'sdata' </code></pre> <p>If I use different variable names for each kernel, that seems to solve the problem. Anyone know why I can't reuse the same variable name? </p>
    singulars
    1. PO
      singulars
      1. PTAnswer
    1. This table or related slice is empty.
    1. PTQuestion
    1. UStalonmies
    1. USmdornfe1
    plurals
    1. PL
      singulars
      1. LTLinked
    1. PL
      singulars
      1. LTLinked
    1. This table or related slice is empty.
    1. PO
      singulars
      1. PTAnswer
    1. VO
      singulars
      1. POCuda reuse of shared memory variable name
      1. This table or related slice is empty.
      1. VTUpMod
    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