 //           + threadIdx.x /*LOCAL_THREAD_X_OFFSET*/;
 //}
 
+// ------------------
+} // namespace __impl
+// ------------------
+
+using namespace __impl;
+
 template <typename T>
 TB_CUDA_GLOBAL void tbCudaFillKernel(T * out, T data, int size)
 {

    if (condition) { printf(__VA_ARGS__); }
#endif

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)


// @formatter:off
tbBOOL tbCudaFill_i(int * out, int data, int count, void * stream)
{ return tbCudaFill(out, data, count, (cudaStream_t)stream) ? TB_TRUE : TB_FALSE; }
tbBOOL tbCudaFill_u(unsigned * out, unsigned data, int count, void * stream)
{ return tbCudaFill(out, data, count, (cudaStream_t)stream) ? TB_TRUE : TB_FALSE; }
tbBOOL tbCudaFill_f(float * out, float data, int count, void * stream)
{ return tbCudaFill(out, data, count, (cudaStream_t)stream) ? TB_TRUE : TB_FALSE; }
tbBOOL tbCudaFill_d(double * out, double data, int count, void * stream)
{ return tbCudaFill(out, data, count, (cudaStream_t)stream) ? TB_TRUE : TB_FALSE; }
// @formatter:on

template <typename T>
