title: OpenCL pass by reference different addres space
body: <p><strong>Short storry:</strong></p>
  
  <p>I have function with pass by reference of output variables</p>
  
  <pre><code>void acum( float2 dW, float4 dFE, float2 *W, float4 *FE )
  </code></pre>
  
  <p>which is supposed increment variables *W, *FE, by dW, dFE if some coditions are fulfilled.
  I want to make this function general - the output varibales can be either local or global.</p>
  
  <pre><code>acum( dW, dFE, &amp;W , &amp;FE  );   // __local acum
  acum( W, FE, &amp;Wout[idOut], &amp;FEout[idOut] );  // __global acum
  </code></pre>
  
  <p>When I try to compile it I got error</p>
  
  <pre><code>error: illegal implicit conversion between two pointers with different address spaces
  </code></pre>
  
  <p>is it possible to make it somehow??? I was thinking if I can use macro instead of function (but I'm not very familier with macros in C).</p>
  
  <p>An other possibility would be probably: </p>
  
  <ol>
  <li>to use function which returns a struct{Wnew, FEnew} </li>
  <li>or (float8)(Wnew,FEnew,0,0) but I don't like it because it makes the code more confusing.</li>
  <li>Naturally I also don't want to just copy the body of "acum" all over the places (like manual inlining it :) )</li>
  </ol>
  
  <p><strong>Backbround (Not necessary to read):</strong></p>
  
  <p>I'm trying program some thermodynamical sampling using OpenCL. Because statistical weight W = exp(-E/kT) can easily overflow float (2^64) for low temperature, I made a composed datatype W = float2(W_digits, W_exponent) and I defined functions to manipulate these numbers "acum". </p>
  
  <p>I try to minimize number of global memory operations, so I let work_items go over Vsurf rather than FEout, becauese I expect that only few points in Vsurf would have considerable statistical weight, so accumulation of FEout would be called only afew times for each workitem. While iteratin over FEout would require sizeof(FEout)*sizeof(Vsurf) memory operations.
  The whole code is here (any recomandations how to make it more efficinet are welcome):</p>
  
  <pre><code>// ===== function :: FF_vdW  - Lenard-Jones Van Der Waals forcefield
  float4 FF_vdW ( float3 R ){
   const float C6    = 1.0;
   const float C12   = 1.0;
   float ir2  = 1.0/ dot( R, R );
   float ir6  = ir2*ir2*ir2;
   float ir12 = ir6*ir6;
   float E6   = C6*ir6;
   float E12  = C12*ir12;
   return (float4)(     
                  ( 6*E6  -  12*E12 ) * ir2    *  R 
                  ,   E12 -     E6
  );}
  
  // ===== function :: FF_spring  - harmonic forcefield
  float4 FF_spring( float3 R){
   const float3 k = (float3)(  1.0, 1.0, 1.0 );
   float3 F = k*R;
   return (float4)(  F,
                     0.5*dot(F,R)
  );}
  
  // ===== function :: EtoW    -   compute statistical weight
  float2 EtoW( float EkT ){
      float Wexp = floor( EkT);
      return (float2)(  exp(EkT - Wexp)
                     ,   Wexp
  ); }
  
  // ===== procedure : addExpInplace    -- acumulate F,E with statistical weight dW
  void acum( float2 dW, float4 dFE, float2 *W, float4 *FE   )
  {
              float dExp = dW.y - (*W).y; // log(dW)-log(W)
              if(dExp&gt;-22){               // e^22 = 2^32 ,    single_float = 2^+64
                 float fac = exp(dExp); 
                 if (dExp&lt;0){             // log(dW)&lt;log(W)
                   dW.x *= fac;
                   (*FE)    += dFE*dW.x;
                   (*W ).x  +=     dW.x;
                 }else{                   // log(dW)&gt;log(W)
                   (*FE)     = dFE  + fac*(*FE); 
                   (*W ).x   = dW.x + fac*(*W).x;
                   (*W ).y   = dW.y;
                 }
              }
  }
  
  // ===== __kernel :: sampler
  __kernel void sampler( 
  __global float  * Vsurf, // in  : surface potential (including vdW)  // may be faster to precomputed endpoints positions like float8
  __global float4 * FEout, // out : Fx,Fy,Fy, E
  __global float2 * Wout,  // out : W_digits, W_exponent
  int3   nV    ,
  float3 dV    ,
  int3   nOut     ,
  int3   iOut0    ,        // shift of Fout in respect to Vsurf
  int3   nCopy    ,        // number of copies of 
  int3   nSample  ,        // dimension of sampling in each dimension around R0   +nSample,-nSample
  float3 RXe0     ,        // postion Xe relative to Tip
  float  EcutSurf ,        
  float  EcutTip  ,        
  float  logWcut  ,        // accumulate only when  log(W) &gt; logWcut
  float  kT                // maximal energy which should be sampled
  ) {
  
      int  id  = get_global_id(0);  // loop over potential grid points
      int  idx = id/nV.x; 
      int3 iV  =  (int3)( idx/nV.y
                        , idx%nV.y
                        , id %nV.x  );
      float V = Vsurf[id];
      float3 RXe = dV*iV;
  
  if (V&lt;EcutSurf){
  // loop over tip position
  for (int iz=0;iz&lt;nOut.z;iz++ ){
  for (int iy=0;iy&lt;nOut.y;iy++ ){
  for (int ix=0;ix&lt;nOut.x;ix++ ){ 
  
      int3   iTip = (int3)( iz, iy, ix );
      float3 Rtip = dV*iTip;
      float4 FE = 0;
      float2 W  = 0;
      // loop over images of potential
      for (int ix=0;ix&lt;nCopy.x;ix++ ){
      for (int iy=0;iy&lt;nCopy.y;iy++ ){
          float3 dR = RXe - Rtip;
          float4 dFE    = FF_vdW( dR );
                 dFE   += FF_spring( dR - RXe0 );
                 dFE.w += V;
  
          if( dFE.w &lt; EcutTip  ){
              float2 dW = EtoW( - FE.w / kT );
              acum( dW, dFE, &amp;W , &amp;FE  );   // __local acum
          }
      } 
      }
  
      if( W.y &gt; logWcut  ){ // accumulate force
          int  idOut = iOut0.x + iOut0.y*nOut.x + iOut0.z*nOut.x*nOut.y;
          acum( W, FE, &amp;Wout[idOut], &amp;FEout[idOut] );  // __global acum
      }
  
  }}}}
  
  }
  </code></pre>
  
  <p>I'm using pyOpenCL on ubuntu 12.04 64bit but I think it has nothink to do with the problem</p>