     if (condition) { printf(__VA_ARGS__); }
 #endif
 
+// ---------------
+namespace __impl {
+// ---------------
+
 template <typename T>
 bool tbCudaGetMaxPotentialBlockSize(int * result_grid_size, int * result_block_size, T func, std::size_t array_count,
                                     std::size_t dynamic_shared_mem_size = 0, int block_size_limit = 0)

//           + threadIdx.x /*LOCAL_THREAD_X_OFFSET*/;
//}

template <typename T>
TB_CUDA_GLOBAL void tbCudaFillKernel(T * out, T data, int size)
{

}

template <typename T, typename StreamType>
bool tbCudaFill(T * out, T data, int count, StreamType stream)
{
    int grid_size(0), block_size(0);
    if (tbCudaGetMaxPotentialBlockSize(&grid_size, &block_size, tbCudaFillKernel<T>, count) == false) {
