 #endif
 }
 
-#ifndef tCudaRawVerbose
-#define tCudaRawVerbose(condition, ...) \
-    if (condition) { \
-        ::sprintf(stdout, __VA_ARGS__)\
-    }
+#ifndef tbCudaRawVerbose
+#define tbCudaRawVerbose(condition, ...) \
+    if (condition) { printf(__VA_ARGS__); }
 #endif
 
 // ---------------

           + threadIdx.x /*LOCAL_THREAD_OFFSET*/;
}

static TB_CUDA_INLINE TB_CUDA_DEVICE int __global_index_1g_2b__()
{
    return (blockIdx.x * (blockDim.x * blockDim.y)) /*GLOBAL_BLOCK_OFFSET*/
           + (threadIdx.y * blockDim.x) /*LOCAL_THREAD_Y_OFFSET*/
           + threadIdx.x /*LOCAL_THREAD_X_OFFSET*/;
}

static TB_CUDA_INLINE TB_CUDA_DEVICE int __global_index_1g_3b__()
{
    return (blockIdx.x * (blockDim.x * blockDim.y * blockDim.z)) /*GLOBAL_BLOCK_OFFSET*/
           + (threadIdx.z * (blockDim.y * blockDim.x)) /*LOCAL_THREAD_X_OFFSET*/
           + (threadIdx.y * blockDim.x) /*LOCAL_THREAD_Y_OFFSET*/
           + threadIdx.x /*LOCAL_THREAD_X_OFFSET*/;
}

template <typename T>
TB_CUDA_GLOBAL void tbCudaAddKernel(T const * v1, T const * v2, T * result, unsigned size)

tbBOOL tbCudaAdd(T const * v1, T const * v2, T * result, unsigned size, StreamType stream)
{
    int grid_size(0), block_size(0);
    if (tbCudaGetMaxPotentialBlockSize(&grid_size, &block_size, tbCudaAddKernel<T>, size) == false) {
        return TB_FALSE;
    }
    assert(grid_size > 0);
    assert(block_size > 0);

    tbCudaAddKernel<T><<<grid_size, block_size, 0, stream>>>(v1, v2, result, size);
    return TB_TRUE;
}

