CUDA Volume Rendering [Part.3 Render(CPU part)]

前面的 part.1 大概介紹了 volume rendering 的概念，也大概列了一下裡面的 function。而 part.2 也已經透過把 main() 掃一遍，把 ray casting 外的東西，大致都講了一下了。而接下來這一部分，主要就來看顯示的部分了∼

在這個程式中，要去把 volume 畫出來的部分，主要就是在 OpenGL 的 display callback function，也就是 display() 這個函式；當每次需要 render 的時候，就會呼叫到這個函式。而整個繪製的流程，大致上如下圖所示：

上圖中，左邊的區域，就是 display() 的大致流程。中間「Render to pixel buffer object」的大框，則是 render() 這個函式在做的事；CUDA 的 kernel 程式 d_Render()，也是在 render() 中被呼叫的。右邊紅色的區域，則是代表 CUDA 的 kernel d_Render() 的流程。

接下來，就先來看 display() 的程式內容：

// use OpenGL to build view matrix 
GLfloat modelView[16]; 
glMatrixMode(GL_MODELVIEW); 
glPushMatrix();
     glLoadIdentity();
     glRotatef(-viewRotation.x, 1.0, 0.0, 0.0);
     glRotatef(-viewRotation.y, 0.0, 1.0, 0.0);
     glTranslatef(-viewTranslation.x, -viewTranslation.y, -viewTranslation.z);
glGetFloatv(GL_MODELVIEW_MATRIX, modelView); 
glPopMatrix();  
invViewMatrix[0] = modelView[0]; invViewMatrix[1] = modelView[4]; invViewMatrix[2] = modelView[8]; invViewMatrix[3] = modelView[12]; 
invViewMatrix[4] = modelView[1]; invViewMatrix[5] = modelView[5]; invViewMatrix[6] = modelView[9]; invViewMatrix[7] = modelView[13];
invViewMatrix[8] = modelView[2]; invViewMatrix[9] = modelView[6]; invViewMatrix[10] = modelView[10]; invViewMatrix[11] = modelView[14]; 
 
render();
  
// display results 
glClear(GL_COLOR_BUFFER_BIT);
  
// draw image from PBO 
glDisable(GL_DEPTH_TEST); 
glRasterPos2i(0, 0); 
glBindBufferARB(GL_PIXEL_UNPACK_BUFFER_ARB, pbo); 
glDrawPixels(width, height, GL_RGBA, GL_UNSIGNED_BYTE, 0); 
glBindBufferARB(GL_PIXEL_UNPACK_BUFFER_ARB, 0);  
glutSwapBuffers(); 
glutReportErrors();

如同前面所述，display() 的內容分成三部分：

1. 處理攝影機的位置矩陣
2. 呼叫 render() 進行繪製
3. 繪製 Pixel Buffer Object

第一部分矩陣計算的部分，基本上這個範例程式在 motion() 中，可以透過滑鼠來控制攝影機（眼睛的位置）的旋轉和平移，並把資料存在 viewRotation 和 viewTranslation 裡；所以在進行 ray casting 前，要先算出來這些影響所形成的矩陣（viewTranslation 在 z 軸預設位移是 -4）。而他的方法是透過 OpenGL 的 model view matrix 來計算，並把結果讀到 modelView 中，然後進行一個轉置，轉換成 kernel 中要用的形式（三乘四的矩陣），也就是 invViewMatrix。

而接下來的第二部分，是把最主要、也就是呼叫 ray casting 的 CUDA kernel 程式，都包到 render() 這個函式；裡面的內容有四步：

1. 將矩陣複製到 device constant memory
2. 將 pixel buffer object pbo 對應到 device memory d_output
3. 呼叫 kernel d_Render()
4. 解除 pbo 和 d_output 的對應

其程式碼如下：

CUDA_SAFE_CALL( cudaMemcpyToSymbol(c_invViewMatrix, invViewMatrix, sizeof(float4)*3) );
  
// map PBO to get CUDA device pointer 
uint *d_output; 
CUDA_SAFE_CALL(cudaGLMapBufferObject((void**)&d_output, pbo));
  
CUDA_SAFE_CALL(cudaMemset(d_output, 0, width*height*4));
  
// call CUDA kernel, writing results to PBO 
d_render<<<gridsize , blockSize>>>(d_output, width, height,
      density, brightness, transferOffset, transferScale);
 
CUT_CHECK_ERROR("kernel failed");
  
CUDA_SAFE_CALL(cudaGLUnmapBufferObject(pbo));

因為在 ray casting 時要用到 invViewMatrix 這個矩陣的內容，而且是要在 kernel 中讀取，所以要先透過 cudaMemcpyToSymbol() 把資料複製到 c_invViewMatrix 這個 constant 變數。而c_invViewMatrix 是一個自訂的 structure，他的定義和宣告如下：

typedef struct {
     float4 m[3]; 
} float3x4;
  
__constant__ float3x4 c_invViewMatrix;  // inverse view matrix
 

接著，就是透過 cudaGLMapBufferObject() 這個函式，來取得之前在 part.2 時所建立的 pixel buffer object pbo 的記憶體位址：d_output；並再透過 cudaMemset() 把他的資料都填 0。

再來就是呼叫 kernel 函式 d_render() 來進行最重要的 ray casting 的動作了！這個 kernel 的參數包含了要寫入結果的 PBO：d_output，以及 pixel buffer object 的寬和高（width, height）；而剩下的 density、brightness、transferOffset、transferScale 這四個變數，主要是用來控制呈現的效果的。關於 kernel 的部分，會在下一部分的 part.4 再來講。

在 kernel 結束後，ray casting 的結果也都會存到 d_output 裡了∼而由於此時已經不會再對 pixel buffer object 做寫入的動作，所以也需要透過 cudaGLUnmapBufferObject()，解除 pbo 和 d_output 之間的對應。

而當 render() 的動作都結束後，就要回到 display() 執行最後一的動作：把 pixel buffer object 畫出來了∼基本上，這邊就全都是 OpenGL 的東西了。主要，就是先透過 glBindBufferARB() 來指定要使用的 buffer object（這邊當然是 pbo），然後再透過 glDrawPixels() 把他畫出來了∼

原始發表：CUDA Volume Rendering [Part.3 Render(CPU part)]