title: cuda calc distance of two points
body: <p>Here I want to calculate the distance of each two points, and decide if they are neighbours. here is my simple code in cuda.</p>
  
  <pre><code>__global__ void calcNeighbors(const DataPoint* points,
    const float doubleRadius, bool* neighbors) {
  
  int tid = threadIdx.x + blockIdx.x * blockDim.x;
  
  float dis = 0.0f;
  while (tid &lt; N) {
     DataPoint p1 = points[tid];
     for (int i=0; i&lt;N; i++) {
         DataPoint p2 = points[i];
         dis = 0;
         dis += (p1.pfDimens[0]-p2.pfDimens[0]) * (p1.pfDimens[0]-p2.pfDimens[0]) +
             (p1.pfDimens[1]-p2.pfDimens[1]) * (p1.pfDimens[1]-p2.pfDimens[1]) +
             (p1.pfDimens[2]-p2.pfDimens[2]) * (p1.pfDimens[2]-p2.pfDimens[2]);
         if (dis &lt;= doubleRadius) {
             neighbors[tid*N+i] = true;
         } else {
             neighbors[tid*N+i] = false;
         }
     }   
  
     tid += blockDim.x * gridDim.x;
  }
  }
  </code></pre>
  
  <p>The DataPoint is a struct is</p>
  
  <pre><code>typedef struct DataPoint {
      float pfDimens[3];
  } DataPoint;
  </code></pre>
  
  <p>so here i want to reduce the time, How can i do? I have tried to use memory coalesing and share memory, but i didn't get a good speed up?</p>
  
  <p>===============use share memory==============</p>
  
  <pre><code>__global__ void calcNeighbors2(const DataPoint* points, 
    const float doubleRadius, bool* neighbors) {
  __shared__ DataPoint sharedpoints[threadsPerBlock];
  
  int start = blockIdx.x * blockDim.x;
  int len = start+threadIdx.x;
  if (len &lt; N) {
      sharedpoints[threadIdx.x] = points[len];
  }
  len = imin(N, blockDim.x + start);
  __syncthreads();
  
  int tid = threadIdx.x;
  float dis;
  while (tid &lt; N) {
      DataPoint p1 = points[tid];
      for (int i=start; i&lt;len; i++) {
         dis = 0;
         dis += (p1.pfDimens[0]-sharedpoints[i-start].pfDimens[0]) * (p1.pfDimens[0]-sharedpoints[i-start].pfDimens[0]) + 
             (p1.pfDimens[1]-sharedpoints[i-start].pfDimens[1]) * (p1.pfDimens[1]-sharedpoints[i-start].pfDimens[1]) + 
             (p1.pfDimens[2]-sharedpoints[i-start].pfDimens[2]) * (p1.pfDimens[2]-sharedpoints[i-start].pfDimens[2]);
         if (dis &lt;= doubleRadius) {
             neighbors[i*N+tid] = true;
         } else {
             neighbors[i*N+tid] = false;
         }
  
      }
  
      tid += blockDim.x;
  }
  }
  </code></pre>
  
  <p>Here i changed the neighbors[tid*N+i] to neighbors[i*N+tid], it give me amlost 8x speed up on Tesla K10.G2.8GB. But when i use share memory to store some points, it is no use? </p>