自己动手，实现在kernel函数中printf()!（转）

在非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）