平行程式設計－CUDA與OpenGL

這次平行的project，我們選擇了一個比較適合cuda的題目，particle system的平行化。
基本上的想法就是使用shader來做render之前的數學計算，讓GPU來加速繪圖過程。
只是因為shader不夠像CUDA，雖然基本上是一樣的，所以我們這次嘗試使用較舊的API，叫做OpenGLCUDA，現在基本上已經棄用的，但是新版的openGL，以及新的nvidia顯示卡依然支援這個語法。




其概念非常簡單，利用繪圖用的buffer，我們可以讓GPU的kernel來run我們的function對buffer中的data做平行化運算，運算完成後才由openGL對buffer中的data做繪圖的動作，現在來看十分的多餘，因為openGL從2010年的3.3版就支援現代的shader語法，讓我們可以簡單地完成這些工作，不過，畢竟學過CUDA，挑戰一下舊語法也可以練習CUDA的一些概念。

首先是openGL，這裡的概念很簡單，建立一個專門記錄particle位置的陣列
叫做vertices，我們先對這個vertices做初始化給好每個點一開始的位置，
因為我們這裡只需要2D的繪圖即可，所以我們不需要用到3維陣列，只要用2維，
當然要做到三維也可以，那麼就會是立體的particle system的動畫。

// Set up vertex data (and buffer(s)) and attribute pointers
GLfloat* vertices = new GLfloat[PARTICLE_SIZE * 8];

接下來我們就是開始用openGL繪圖，這裡我們還不用更新每個位置，只是單純化出particle而已，這裡要特別看一下VBO，因為接下來這個部分會相當重要。

GLuint VBO;
glGenBuffers(1, &VBO);
glBindBuffer(GL_ARRAY_BUFFER, VBO);
glBufferData(GL_ARRAY_BUFFER, PARTICLE_SIZE * 8 * sizeof(GLfloat), 0, GL_DYNAMIC_DRAW);

glGenBuffers
這個語法就是現代openGL關鍵的東西，應該說接下來的GL都是用相同的概念在畫圖的
以前的GL是非常簡單的一個口令一個動作，
也就是GLDRAW、GLRotate....像是用輸入指令的方式對著螢幕一筆一筆地畫出圖案。
當然這是很浪費時間的，不管對電腦還是對人來說。
於是新的GL開始使用陣列方式去儲存每一個點，每一個點的位置排好，接下來一樣GLDrawArray，GL會一口氣把目標陣列中的所有點畫到螢幕上。
而這個陣列就是要用glGenBuffers告訴openGL，接下來的GLDrawArray的目標陣列是誰。
openGL的溝通就是用VBO這樣的int來知道是哪個array，
所以我們生成一個VBO int等等會拿來代表目標陣列的整數，
再用glBindBuffer(GL_ARRAY_BUFFER, VBO);代表接下來的GLDrawArray的目標陣列是要找哪個代表整數
所以接下來的glBufferData就會擁有這個VBO的數字，以後VBO在我們的openGL程式中就可以代表這個array，所以glBufferData的第二個參數是要多大的陣列，第三個參數則是這個陣列要複製哪個來源的內容，0就是NULL代表不用複製。

看到這應該會有些疑問，那所謂的VAO呢？
有了VBO應該就可以做完所有的事了不是嗎？
的確，但是到了現在的繪圖，我們希望openGL可以做到更多的事，
所以我們嘗試對VBO代表的目標陣列中的data可以去修改，
當然如果每次畫圖都對陣列做修改，想必花費的時間一定更多，
但是我們如果要做動畫甚至是遊戲，怎麼可能不修改data資料呢？
所以GL也讓我們可以對目標陣列中的data做修改，為了加速修改過程的時間，
我們將會只用GPU來做這個運算，那麼就必須要用VAO
VAO會去指定目標array中的特定位置的所有data，並給他一個VAO數字代表他。
所以看出來了，VAO就是把VBO中的資料取出想修改的部分拿到GPU裡運算後在整個畫出
所以語法當然就會跟VBO息息相關了，畢竟是要代表VBO想修改的地方，

GLuint VBO, VAO;
glGenVertexArrays(1, &VAO);
glGenBuffers(1, &VBO);
// Bind the Vertex Array Object first, then bind and set vertex buffer(s) and attribute pointer(s).
glBindVertexArray(VAO);
glBindBuffer(GL_ARRAY_BUFFER, VBO);
glBufferData(GL_ARRAY_BUFFER, PARTICLE_SIZE * 8 * sizeof(GLfloat), 0, GL_DYNAMIC_DRAW);

glVertexAttribPointer(0, 3, GL_FLOAT, GL_FALSE, 6 * sizeof(GLfloat), (GLvoid*)0);
glEnableVertexAttribArray(0);

glVertexAttribPointer(1, 3, GL_FLOAT, GL_FALSE, 6 * sizeof(GLfloat), (GLvoid*)(3 * sizeof(GLfloat)));
glEnableVertexAttribArray(1);
glBindVertexArray(0); // Unbind VAO

上色部分就是新增的VAO語法，可以看到一樣也是使用glBindVertexArray來代表接下的參考陣列是誰，glEnableVertexAttribArray就是指定說這些位置上的data的編號是多少，等等GPU運算就可以用這個編號來統一對這些點做計算。
這裡分成兩個編號，一個是0,1,2,6,7,8,12,13,14.......編號為0
另外一個就是3,4,5,9,10,11,15,16,17......編號為1

寫完編號之後要怎麼傳到GPU理運作呢？
openGL把這個平行化的function命名為shader
shader裡面長這樣

#version 330 core
layout (location = 0) in vec3 position;
layout (location = 2) in vec3 ourcolor;
out vec3 color
void main()
{
    gl_Position = vec4(position, 1.0f);
    color = ourcolor;
}

因此只要在GLDrawArray之前use shade就可以針對每個要處理的data先用shader運算之後再交由openGL一口氣畫出來。


glBufferData(GL_ARRAY_BUFFER, PARTICLE_SIZE * 8 * sizeof(GLfloat), 0, GL_DYNAMIC_DRAW);

glBufferData(GL_ARRAY_BUFFER, PARTICLE_SIZE * 8 * sizeof(GLfloat), ParticleArray, GL_DYNAMIC_DRAW);

這兩個程式碼的差別只是在於第3個參數，我們是否需要複製一次資料內容到buffer array裡面，而資料內容的來源用一個陣列的指標告知在哪裡，0就是代表NULL，不用複製任何資料到buffer data裡。

基本的openGL都設定完了之後，就是CUDA程式碼的部分。
首先是註冊VBO，讓CUDA可以知道VBO是誰，所以我們註冊VBO到CUDA裡
cudaGraphicsResource_t resource;
cudaGraphicsGLRegisterBuffer(&resource, VBO, cudaGraphicsMapFlagsNone);

如此一來VBO這個buffer object就會被存到resource中，以後CUDA想要更動VBO就可以直接call resource來取得VBO中的資料內容。

接下來是CUDA的資料輸入的話，其type一定要是void **，所以我們再多做一步，把resource這個pointer轉成 void **的型態，這個動作CUDA也有提供API，這個動作叫做「map resource」

float* devPtr;
size_t size;
cudaGraphicsMapResources(1, &resource, NULL);
cudaGraphicsResourceGetMappedPointer((void**)&devPtr, &size, resource);

如此一來devPtr就可以拿來當作資料輸入的參數了。
接下來就是最基本的cuda memory copy的動作

cudaMemcpy(devPtr, vertices, PARTICLE_SIZE * 8 * sizeof(float), cudaMemcpyHostToDevice);

接下來launch kernel。

dim3 threadsPerBlock(1024);
dim3 numBlocks(PARTICLE_SIZE / 1024 + 1);
updateParticles << < numBlocks, threadsPerBlock >> > (devPtr, mousePos, LMB, dt);

接下來別忘記unmap cuda。

cudaGraphicsUnmapResources(1, &resource, 0); // give access authority of vbo1 back to openGL
cudaGraphicsUnregisterResource(resource); // unregiste the resource  

這個非常重要，必須要unmap cuda，GL才會畫圖，否則GL無法使用buffer來畫圖。
當然unmap後，下次要lunch kernel就再次map buffer即可。