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

//           + 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) {
