在非EMU模式下，CUDA kernel函数中不能调用任何标准库函数，自然也就不能调用printf。这给我们调试CUDA程序带来了极大的不便。为此，最近nVidia已经专门开发了一个cuprintf()函数，不过目前还在内部测试中，普通用户暂时还无法看到。
其实，我们利用CUDA现有的功能自行开发一个类似的函数也并不是很困难。自己动手，丰衣足食。本文中的给出的小程序就是这样一个简单的实现。这是一个完整的测试程序，可以直接在VS2008上运行。在真正的device中printf()和在EMU模式下printf()感觉是不一样的，大家可以尝试一下。
使用方法和局限性：
1、只要首先在host端调用device_printf_init()初始化，然后就可以在kernel函数中像printf()一样调用device_printf()了。
2、由于device函数一般不支持可变参数，为简单起见，本文实现的device_printf只支持1个或2个参数，暂时只支持'%d', '%u', '%f', '%c'这几种描述符。以下用法都是允许的：
device_printf("1+1=%d/n", 2);
device_printf("A+B=%c/n", 'C');
device_printf("1/2=%f/n", 1/2.0);
device_printf("unsigned(-1)=%u/n", -1);
device_printf("OK, test passed./n");
3、输出结果会自动flush(), 默认flush()时间是200ms，可以通过修改宏参数PRINTF_FLUSH_INTERVAL调整。
4、打印缓冲区默认为4096 (PRINTF_BUFFSER_SIZE)，如果打印太快(200ms内超过4096字节)会发生溢出。溢出的后果是造成部分打印结果丢失。其他没什么副作用。
由于程序写得匆忙，还没有仔细检查。欢迎报告bug。
-
C/C++ code
-
/* * device_printf.h - my own implemetation of 'cuprintf', enable limited printf() from device * v0.1 by cuda2010 @ csdn, Feb 21, 2010 */ #ifndef DEVICE_PRINFT_H #define DEVICE_PRINTF_H #include " cutil_inline.h " #include " windows.h " #pragma comment(lib, "winmm.lib") using namespace std; #define PRINTF_BUF_SIZE 4096 #define PRINTF_FLUSH_INTERVAL 200 struct BUFFER { char data[PRINTF_BUF_SIZE]; unsigned int ptr; } * host_buffer; __device__ struct BUFFER * device_buffer; static void CALLBACK device_printf_flush(UINT uid, UINT umsg, DWORD dwusr, DWORD dw1, DWORD dw2) { static int old_ptr; int ptr = (host_buffer -> ptr) % PRINTF_BUF_SIZE; if (old_ptr <= ptr) { for ( int i = old_ptr; i < ptr; i ++ ) { printf( " %c " , host_buffer -> data[i]); } } else { for ( int i = old_ptr; i < PRINTF_BUF_SIZE; i ++ ) printf( " %c " , host_buffer -> data[i]); for ( int i = 0 ; i < ptr; i ++ ) printf( " %c " , host_buffer -> data[i]); } old_ptr = ptr; } __global__ void set_device_buffer( void * ptr) { device_buffer = ( struct BUFFER * )ptr; } void device_printf_init() { static int inited = 0 ; void * ptr; if (inited) return ; timeSetEvent(PRINTF_FLUSH_INTERVAL, 0 , device_printf_flush, 0 , (UINT)TIME_PERIODIC); cutilSafeCall(cudaSetDeviceFlags(cudaDeviceMapHost)); cutilSafeCall(cudaHostAlloc(( void ** ) & host_buffer, sizeof ( struct BUFFER), cudaHostAllocMapped)); memset(host_buffer, 0 , sizeof ( struct BUFFER)); cutilSafeCall(cudaHostGetDevicePointer(( void ** ) & ptr, host_buffer, 0 )); set_device_buffer <<< 1 , 1 >>> (ptr); cudaThreadSynchronize(); inited = 1 ; } __device__ int sprintf_int( char * s, int v) { int len = 0 ; if (v == 0 ) { s[ 0 ] = ' 0 ' ; return (len = 1 ); } if (v < 0 ) { s[len ++ ] = ' - ' ; v =- v; } unsigned int base ; for ( base = 1000000000 ; base > 0 ; base /= 10 ) if ((unsigned int )v >= base ) break ; while ( base > 0 ) { s[len ++ ] = ' 0 ' + (v / base ) % 10 ; base /= 10 ; } return len; } __device__ int sprintf_uint( char * s, unsigned int v) { int len = 0 ; if (v == 0 ) { s[ 0 ] = ' 0 ' ; return (len = 1 ); } unsigned int base ; for ( base = 1000000000 ; base > 0 ; base /= 10 ) if (v >= base ) break ; while ( base > 0 ) { s[len ++ ] = ' 0 ' + (v / base ) % 10 ; base /= 10 ; } return len; } __device__ int sprintf_float( char * s, double v) { int len = 0 ; if (isinf(v)) { s[len ++ ] = ' I ' ; s[len ++ ] = ' N ' ; s[len ++ ] = ' F ' ; return len; } if (isnan(v)) { s[len ++ ] = ' N ' ; s[len ++ ] = ' A ' ; s[len ++ ] = ' N ' ; return len; } if (v < 0 ) { s[len ++ ] = ' - ' ; v =- v; } double e = log10(v); int a, b; if (isinf(e)) e = 0 ; b = ( int )e; if (b < 0 ) b -- ; a = ( int )(v / pow( 10.0 ,b) * 1000000 + 0.5 ); s[len ++ ] = ' 0 ' + (a / 1000000 ) % 10 ; s[len ++ ] = ' . ' ; s[len ++ ] = ' 0 ' + (a / 100000 ) % 10 ; s[len ++ ] = ' 0 ' + (a / 10000 ) % 10 ; s[len ++ ] = ' 0 ' + (a / 1000 ) % 10 ; s[len ++ ] = ' 0 ' + (a / 100 ) % 10 ; s[len ++ ] = ' 0 ' + (a / 10 ) % 10 ; s[len ++ ] = ' 0 ' + (a) % 10 ; s[len ++ ] = ' E ' ; s[len ++ ] = (b < 0 ) ? ' - ' : ' + ' ; b = abs(b); s[len ++ ] = ' 0 ' + (b / 100 ) % 10 ; s[len ++ ] = ' 0 ' + (b / 10 ) % 10 ; s[len ++ ] = ' 0 ' + (b) % 10 ; return len; } __device__ int sprintf_str( char * s, char * s2, int maxlen) { int i; for (i = 0 ; s2[i] && i < maxlen; i ++ ) s[i] = s2[i]; return i; } __device__ int do_device_printf( char * s, int len) { unsigned int from = atomicAdd( & (device_buffer -> ptr), len); for ( int i = 0 ; i < len; i ++ ) device_buffer -> data[(from + i) & (PRINTF_BUF_SIZE - 1 )] = s[i]; return len; } template < class TYPE > __device__ int device_printf(TYPE fmt) { char * s = ( char * )fmt; int len; for (len = 0 ; s[len] && len < 255 ; len ++ ); do_device_printf(s, len); return len; } template < class TYPE > __device__ int device_printf( char * fmt, TYPE a) { char buf[ 300 ]; int i, j; for (i = j = 0 ; fmt[i]; i ++ ) { if (j >= 255 ) break ; if (fmt[i] != ' % ' ) { buf[j ++ ] = fmt[i]; continue ; } if (fmt[i + 1 ] == ' d ' ) { j += sprintf_int(buf + j, a); } else if (fmt[i + 1 ] == ' u ' ) { j += sprintf_uint(buf + j, a); } else if (fmt[i + 1 ] == ' c ' ) { buf[j ++ ] = a; } else if (fmt[i + 1 ] == ' f ' ) { j += sprintf_float(buf + j, ( double )a); } i += 2 ; j += sprintf_str(buf + j, fmt + i, 255 - j); break ; } if (j > 255 ) j = 255 ; buf[j] = 0 ; return do_device_printf(buf, j); } #endif // end of DEVICE_PRINFT_H // main.c __global__ void cuda1() { int inx = blockIdx.x * blockDim.x + threadIdx.x; float sum = 0 ; device_printf( " Thread %d starts/n " , inx); for ( int i = 1 ; i < 20000000 ; i ++ ) { sum += ( float )i * i / 10.0f ; if (i % 1000000 == 0 ) { device_printf( " Thread %d:/n " , inx); device_printf( " i=%d/n " , i); device_printf( " sum=%f/n " , sum); } } device_printf( " Thread (%c) is finished/n " , inx + ' A ' ); } int main() { device_printf_init(); cuda1 <<< 4 , 4 >>> (); cudaThreadSynchronize(); }
-
（转自：http://topic.csdn.net/u/20100221/12/85d9155c-f7d3-4edf-a3f6-3af3d0bb6835.html）
53
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
