8.4 基于图形互操作性的热传导

那么，前面介绍的这些内容的要点在哪里？如果观察在前面动画示例中使用的CPUAnimBitmap结构，我们会发现它的工作原理与8.2节中的绘制代码非常相似。

CPUAnimBitmap与前一个示例的关键差异在于对glDrawPixels()的调用。

glDrawPixels( bitmap->x,

bitmap->y,
GL_RGBA,
GL_UNSIGNED_BYTE,
bitmap->pixels);

在本章的第一个示例中，你可能已经看到了对glDrawPixels()的调用，其中最后的参数是一个缓冲区指针。如果你之前没有注意，那么现在就看到了。在调用CPUAnimBitmap的Draw()函数时，将把bitmap->pixels中的CPU缓存复制到GPU中以便进行绘制。为了实现这个操作，CPU需要停止正在执行的工作，并将每一帧都复制到GPU。这需要在CPU与GPU之间进行同步。此外，在PCI Express总线上启动和完成传输还会导致额外的延迟。由于在glDrawPixels()的调用中最后一个参数是主机指针，这就意味着，在通过CUDA C核函数生成一帧图像数据后，需要通过CUDAMemcpy()将该帧从GPU复制到CPU。

void generate_frame(DataBlock \*d，int ticks）{ dim3 grids(DIM/16, DIM/16); dim3 threads(16,16); kernel<<<grids,threads>>>(d->dev bitmap，ticks); HANDLE_ERROR( CUDAMemcpy(d-> bitmap->get_ptr(), d->dev bitmap, d-> bitmap->image_size(), CUDAMemcpyDeviceToHost)）;   
}

总的来说，在最初的GPU波纹应用程序中存在着许多可以改进的地方。这个程序使用CUDA来计算在每一帧绘制的图像，但在计算完成后，我们将缓冲区复制到CPU，然后又将缓冲区复制回GPU以显示。可以看到，在主机与设备之间存在着不必要的数据迁移，使得程序无法实现最优性能。接下来，我们将重新回顾计算密集的动画应用程序，并通过图形互操作来实现绘制操作从而实现性能提升。

如果回顾第7章的热模拟应用程序，你会发现其中同样使用了CPUAnimBitmap来显示模拟计算的输出结果。我们将把这个应用程序修改为使用新实现的GPUAnimBitmap结构，并观察性能将发生怎样的变化。与波纹示例一样，GPUAnimBitmap可以很容易地替代CPUAnimBitmap，只需将unsigned char改为uchar4即可。因此，我们修改动画函数的原型从而适应这种在数据类型上的变化。

void anim_gpu( uchar4* outputBitmap, DataBlock *d, int ticks) {
    HANDLE_ERROR( CUDAEventRecord(d->start, 0));
    dim3 blocks(DIM/16, DIM/16);
    dim3 threads(16, 16);

// 由于tex是全局的并且有界的，我们必须使用一个标志来表示

// 在每次迭代中哪个是输入以及哪个是输出
volatile bool dstOut = true;
for (int i=0; i<90; i++) {
float *in, *out;
if (dstOut) {
in = d->dev_inSrc;
out = d->dev_outSrc;
} else {
out = d->dev_inSrc;
in = d->dev_outSrc;
}
copy_const_kernel<<blocks, threads>>(in);
blend_kernel<<blocks, threads>>(out, dstOut);
dstOut = !dstOut;
}
float_to_color<<blocks, threads>>(outputBitmap,
d->dev_inSrc);
HANDLE_ERROR(udaEventRecord(d->stop, 0));
HANDLE_ERROR(udaEventSynchronize(d->stop));
float elapsedTime;
HANDLE_ERROR(udaEventElapsedTime(& elapsedTime,
d->start, d->stop));
d->totalTime += elapsedTime;
++d->frames;
printf("Average Time per frame: %3.1f ms\n", d->totalTime/d->frames);
}

由于核函数float_to_color()是唯一一个使用outputBitmap的函数，因此在将数据类型修改为uchar4后，只需对这个函数进行修改。在第7章中，只认为这个函数是一段辅助代码，在这里将仍然如此。但是，我们重载了这个函数，并在book.h中同时包含了unsigned char和uchar4的版本。你将注意到，这些函数间的差异与GPU波纹示例中不同版本kernel()之间的差异是类似的。为了简单，在这里省略了float_to_color()核函数的大部分代码，如果需要进一步观察，可以参考book.h。

__global__ void float_to_color( unsigned char *optr, const float *outSrc) {

// 将浮点值转换为4个颜色值

optr[offset*4 + 0] = value(m1, m2, h+120);
optr[offset*4 + 1] = value(m1, m2, h);
optr[offset*4 + 2] = value(m1, m2, h -120);

optr[offset*4 + 3] = 255;  
}  
__global__void float_to_color(uchar4 *optr, const float *outSrc) {

//将浮点值转换为4个颜色值

optr[offset].x = value(m1, m2, h+120);
optr[offset].y = value(m1, m2, h);
optr[offset].z = value(m1, m2, h-120);
optr[offset].w = 255;

除了这些修改外，另一个主要的差异在于将CPUAnimBitmap修改为GPUAnimBitmap以便执行动画操作。

int main(void) { DataBlock data; GPUAnimBitmap bitmap(DIM, DIM, &data); data.totalTime = 0; dataFrames = 0; HANDLE_ERROR(udaEventCreate( &data.start)); HANDLE_ERROR(udaEventCreate( &data.stop)); int imageSize = bitmap.image_size(); //假设float类型的大小为4字符（即rgb） HANDLE_ERROR(udaMalloc( (void**)&data.dev_inSrc, imageSize)); HANDLE_ERROR(udaMalloc( (void**)&data.dev_outSrc, imageSize)); HANDLE_ERROR(udaMalloc( (void**)&data.dev_constSrc, imageSize)); HANDLE_ERROR(udaBindTexture( NULL, texConstSrc, data.dev_constSrc, imageSize)); HANDLE_ERROR(udaBindTexture( NULL, texIn, data.dev_inSrc, imageSize)); HANDLE_ERROR(udaBindTexture( NULL, texOut, data.dev_outSrc, imageSize));

//初始化常量数据

float *temp = (float*)malloc( imageSize );   
for (int i=0; i< DIM*DIM; i++) { temp[i] = 0; int x = i % DIM; int y = i / DIM; if ((x>300) && (x<600) && (y>310) && (y<601)) temp[i] = MAX_TEMP; } temp[ DIM*100+100] = (MAX_TEMP + MIN_TEMP)/2; temp[ DIM*700+100] = MIN_TEMP; temp[ DIM*300+300] = MIN_TEMP; temp[ DIM*200+700] = MIN_TEMP; for (int y=800; y<900; y++) { for (int x=400; x<500; x++) { temp[x+y*DIM] = MIN_TEMP; } }   
}   
HANDLE_ERROR(udaMemcpy( data.dev_constSrc, temp, imageSize,udaMemcpyHostToDevice ) );

//初始化输入数据

for (int y=800; y<DIM; y++) { for (int x=0; x<200; x++) { temp[x+y*DIM] = MAX_TEMP; } }   
HANDLE_ERROR(udaMemcpy(data.dev_inSrc, temp, imageSize,udaMemcpyHostToDevice) ); free(temp);   
bitmap anim_and_exit((void*)(uchar4*,void*,int))anim_gpu, (void*)(void*))anim_exit);

虽然在改进前后的热模拟程序中还存在一些差异，但它们并不重要，在这里也就不进一步讨论了。最重要的是回答这个问题：当我们把程序完全迁移到GPU上后，程序的性能将发生怎样的改变？现在，程序无需将每一帧都复制回主机进行显示，因此比之前的情况更为乐观。

那么，基于图形互操作行来实现图像绘制将获得哪种程度的性能提升？之前，在基于 GeForce GTX 285 的测试机器上，热传导示例大概每帧消耗 25.3 毫秒。在将应用程序修改为使用图形互操作性后，降低到了每帧 21.6 毫秒。结果就是绘制操作快了 $15\%$ ，并且在每次显示一

帧图像时不再需要主机的介入。这是个还不错的结果！

8.4_基于图形互操作性的热传导

8.4 基于图形互操作性的热传导