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 的時候,就會呼叫到這個函式。而整個繪製的流程,大致上如下圖所示:

render1

上圖中,左邊的區域,就是 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. 解除 pbod_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);而剩下的 densitybrightnesstransferOffsettransferScale 這四個變數,主要是用來控制呈現的效果的。關於 kernel 的部分,會在下一部分的 part.4 再來講。

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

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


  • Part.1 簡介
  • Part.2 程式流程
  • Part.3 Render(CPU part)
  • Part.4 Render(GPU part)

    nVidia CUDA 相關文章目錄

  • 對「CUDA Volume Rendering [Part.3 Render(CPU part)]」的想法

    發表迴響

    在下方填入你的資料或按右方圖示以社群網站登入:

    WordPress.com 標誌

    您的留言將使用 WordPress.com 帳號。 登出 /  變更 )

    Google photo

    您的留言將使用 Google 帳號。 登出 /  變更 )

    Twitter picture

    您的留言將使用 Twitter 帳號。 登出 /  變更 )

    Facebook照片

    您的留言將使用 Facebook 帳號。 登出 /  變更 )

    連結到 %s

    This site uses Akismet to reduce spam. Learn how your comment data is processed.