title: Cuda Shared Memory array variable
body: <p>I am trying to declare a variable for matrix multiplication as follows:</p>
  
  <pre><code>__shared__ float As[BLOCK_SIZE][BLOCK_SIZE];
  </code></pre>
  
  <p>I am trying to make it so the user could input the size of the matrix to calculate, however that would mean changing the BLOCK_SIZE. I changed it but I am getting a compiler error:"error: constant value is not known". I've looked into it and it's similar to this <a href="https://stackoverflow.com/questions/3220553/define-variable-size-on-array-in-local-memory-using-cuda">thread</a>. So I tried:</p>
  
  <pre><code>__shared__ int buf [];
  </code></pre>
  
  <p>But then I get: "error: incomplete type is not allowed"</p>
  
  <p>Thanks,
  Dan
  Update with code(pretty much followed <a href="http://gpgpu-computing4.blogspot.com/2009/09/matrix-multiplication-3.html" rel="nofollow noreferrer">this guide</a> and the staring out with cuda guide):
  The block size is passed in by asking the user of the size of the matrix. They enter the x and y. Block size is only x and right now it has to accept the same size as x and y.</p>
  
  <pre><code>__global__ void matrixMul( float* C, float* A, float* B, int wA, int wB,size_t block_size)
  {
      // Block index
      int bx = blockIdx.x;
      int by = blockIdx.y;
  
      // Thread index
      int tx = threadIdx.x;
      int ty = threadIdx.y;
  
      // Index of the first sub-matrix of A processed 
      // by the block
      int aBegin = wA * block_size * by;
  
      // Index of the last sub-matrix of A processed 
      // by the block
      int aEnd   = aBegin + wA - 1;
  
      // Step size used to iterate through the 
      // sub-matrices of A
      int aStep  = block_size;
  
      // Index of the first sub-matrix of B processed 
      // by the block
      int bBegin = block_size * bx;
  
      // Step size used to iterate through the 
      // sub-matrices of B
      int bStep  = block_size * wB;
      float Csub=0;
      // Loop over all the sub-matrices of A and B
      // required to compute the block sub-matrix
      for (int a = aBegin, b = bBegin; a &lt;= aEnd; a += aStep, b += bStep) 
      {
          // Declaration of the shared memory array As 
          // used to store the sub-matrix of A
  
          extern __shared__ float As[];
  
          // Declaration of the shared memory array Bs 
          // used to store the sub-matrix of B
          extern __shared__ float Bs[];
          extern __shared__ float smem[];
  
          // Load the matrices from global memory
          // to shared memory; each thread loads
          // one element of each matrix
          smem[ty*block_size+tx] = A[a + wA * ty + tx];
          //cuPrintf("\n\nWhat are the memory locations?\n");
          //cuPrintf("The shared memory(A) is: %.2f\n",smem[ty*block_size+tx]);
          smem[block_size*block_size+ty*block_size+tx]  = B[b + wB * ty + tx];
          //cuPrintf("The shared memory(B) is: %.2f\n",smem[block_size*block_size+ty*block_size+tx]);
          // Synchronize to make sure the matrices 
          // are loaded
          __syncthreads();
  
          // Multiply the two matrices together;
          // each thread computes one element
          // of the block sub-matrix
          for (int k = 0; k &lt; block_size; ++k)
          {
  
              Csub += smem[ty*block_size+k] * smem[block_size*block_size+k*block_size+tx] ;
              //cuPrintf("Csub is currently: %.2f\n",Csub);
          }
          //cuPrintf("\n\n\n");
          // Synchronize to make sure that the preceding
          // computation is done before loading two new
          // sub-matrices of A and B in the next iteration
          //cuPrintf("the results are csub: %.2f\n",Csub);
          __syncthreads();
      }
      // Write the block sub-matrix to device memory;
      // each thread writes one element
      int c = wB * block_size * by + block_size * bx;
      C[c + wB * ty + tx] = Csub;
  
  
  }
  </code></pre>