         const float bandwidth_hz, const float time_int_sec,
         const float gha0_rad, const float dec0_rad, float4c* d_vis)
 {
+#if __CUDA_ARCH__ >= 300
     /* Common values per thread block. */
     __shared__ float uu, vv, du_dt, dv_dt, dw_dt;
     __shared__ const float4c* __restrict__ station_i;

        const float bandwidth_hz, const float time_int_sec,
        const float gha0_rad, const float dec0_rad, float4c* d_vis)
{
    // This kernel evaluates baselines p-q for a single station p and a range
    // of stations q, determined by the BLOCK_SIZE


            d_vis[ipq].d.y += b_sum[threadIdx.x].d.y;
        }
    }
#endif /* __CUDA_ARCH__ >= 300 */
}

void correlate_warpshuffle_blocked_f(int num_sources, int num_stations, const float4c* d_Jones,
