OpenGL與CUDA互操作方式總結

一、介紹

CUDA是Nvidia推出的一個通用GPU計算平台，對於提升並行任務的效率非常有幫助。本人主管的項目中采用了OpenGL做圖像渲染，但是在數據處理方面比較慢，導致幀率一直上不來。於是就嘗試把計算工作分解成小的任務，使用核函數在CUDA中加速計算。對於CUDA和OpenGL如何交互以前從來沒有接觸過，這次在實施時趟了不少的坑。在這里記錄下OpenGL與CUDA的互操作的兩種方式。

二、基本操作流程

OpenGL與CUDA互操作可以分成兩種，一種是OpenGL將Buffer對象注冊到CUDA中去，供CUDA讀寫操作，然后再在OpenGL中使用。一般這種情況下注冊的是VBO和PBO，VBO一般用於存儲頂點坐標、索引等數據；PBO則一般用於存儲圖像數據，因此稱作Pixel Buffer Object。另一種是OpenGL將Texture對象注冊到CUDA中去，經CUDA處理后得到紋理內容，然后在OpenGL中渲染出來。不過不管是哪一種互操作類型，其操作流程是一致的：

• 在OpenGL里面初始化Buffer Object
• 在CUDA中注冊OpenGL中的Buffer Object
• CUDA鎖定資源，獲取操作資源的指針，在CUDA核函數中進行處理
• CUDA釋放資源，在OpenGL中使用Buffer Object

下面就以代碼為例，講講兩種方式的異同：

（1）OpenGL PBO/VBO在CUDA中的使用

// 初始化Buffer Object
//vertex array object
glGenVertexArrays(1, &this->VAO);
//Create vertex buffer object
glGenBuffers(2, this->VBO);
//Create Element Buffer Objects
glGenBuffers(1, &this->EBO);

//Bind the Vertex Array Object first, then bind and set vertex buffer(s) and attribute pointer(s).
glBindVertexArray(this->VAO);

// 綁定VBO后即在CUDA中注冊Buffer Object
glBindBuffer(GL_ARRAY_BUFFER, this->VBO[0]);
glBufferData(GL_ARRAY_BUFFER, sizeof(*this->malla)*this->numPoints, this->malla, GL_DYNAMIC_COPY);
cudaGraphicsGLRegisterBuffer(&this->cudaResourceBuf[0], this->VBO[0], cudaGraphicsRegisterFlagsNone);

glBindBuffer(GL_ARRAY_BUFFER, this->VBO[1]);
glBufferData(GL_ARRAY_BUFFER, sizeof(*this->malla)*this->numPoints, this->malla, GL_DYNAMIC_COPY);
cudaGraphicsGLRegisterBuffer(&this->cudaResourceBuf[1], this->VBO[1], cudaGraphicsRegisterFlagsNone);

// 在CUDA中映射資源，鎖定資源
cudaGraphicsMapResources(1, &this->cudaResourceBuf[0], 0);
cudaGraphicsMapResources(1, &this->cudaResourceBuf[1], 0);

point *devicePoints1;
point *devicePoints2;
size_t size = sizeof(*this->malla)*this->numPoints;
// 獲取操作資源的指針，以便在CUDA核函數中使用
cudaGraphicsResourceGetMappedPointer((void **)&devicePoints1, &size, this->cudaResourceBuf[0]);
cudaGraphicsResourceGetMappedPointer((void **)&devicePoints2, &size, this->cudaResourceBuf[1]);
// execute kernel
dim3 dimGrid(20, 20, 1);
dim3 dimBlock(this->X/dimGrid.x, this->Y/dimGrid.y, 1);
modifyVertices<<<dimGrid, dimBlock>>>(devicePoints1, devicePoints2,this->X, this->Y);
modifyVertices<<<dimGrid, dimBlock>>>(devicePoints2, devicePoints1,this->X, this->Y);

// 處理完了即可解除資源鎖定，OpenGL可以開始利用處理結果了。
// 注意在CUDA處理過程中，OpenGL如果訪問這些鎖定的資源會出錯。
cudaGraphicsUnmapResources(1, &this->cudaResourceBuf[0], 0);
cudaGraphicsUnmapResources(1, &this->cudaResourceBuf[1], 0);

　　值得注意的是，由於這里綁定的是VBO，屬於Buffer對象，因此調用的CUDA API是這兩個：

cudaGraphicsGLRegisterBuffer();
cudaGraphicsResourceGetMappedPointer();

（2）OpenGL Texture在CUDA中的使用

// 初始化兩個Texture並綁定
cudaGraphicsResource_t cudaResources[2];
GLuint textureID[2];
glEnable(GL_TEXTURE_2D);
glGenTextures(2, textureID);
glBindTexture(GL_TEXTURE_2D, textureID[0]);
glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST);
glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_NEAREST);
glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA, 1000, 1000, 0, GL_RGBA, GL_UNSIGNED_BYTE, NULL);

glBindTexture(GL_TEXTURE_2D, textureID[1]);
glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST);
glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_NEAREST);
glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA, 1000, 1000, 0, GL_RGBA, GL_UNSIGNED_BYTE, NULL);
// 在CUDA中注冊這兩個Texture
cudaError_t err = cudaGraphicsGLRegisterImage(&cudaResources[0], textureID[0], GL_TEXTURE_2D, cudaGraphicsRegisterFlagsWriteDiscard);
if (err != cudaSuccess)
{
    std::cout << "cudaGraphicsGLRegisterImage: " << err << "Line: " << __LINE__;
    return -1;
}
err = cudaGraphicsGLRegisterImage(&cudaResources[1], textureID[1], GL_TEXTURE_2D, cudaGraphicsRegisterFlagsWriteDiscard);
if (err != cudaSuccess)
{
    std::cout << "cudaGraphicsGLRegisterImage: " << err << "Line: " << __LINE__;
    return -1;
}
// 在CUDA中鎖定資源，獲得操作Texture的指針，這里是CudaArray*類型
cudaError_t err = cudaGraphicsMapResources(2, cudaResource, 0);	
err = cudaGraphicsSubResourceGetMappedArray(&this->cuArrayL, cudaResource[0], 0, 0);
err = cudaGraphicsSubResourceGetMappedArray(&this->cuArrayR, cudaResource[1], 0, 0);

// 數據拷貝至CudaArray。這里因為得到的是CudaArray，處理時不方便操作，於是先在設備內存中
// 分配緩沖區處理，處理完后再把結果存到CudaArray中，僅僅是GPU內存中的操作。
cudaMemcpyToArray(cuArrayL, 0, 0, pHostDataL, imgWidth*imgHeight * sizeof(uchar4), cudaMemcpyDeviceToDevice);
cudaMemcpyToArray(cuArrayR, 0, 0, pHostDataR, imgWidth*imgHeight * sizeof(uchar4), cudaMemcpyDeviceToDevice);
// 處理完后即解除資源鎖定，OpenGL可以利用得到的Texture對象進行紋理貼圖操作了。
cudaGraphicsUnmapResources(1, &cudaResource[0], 0);
cudaGraphicsUnmapResources(1, &cudaResource[1], 0);

 

      注意這里因為使用的是Texture對象，因此使用了不同的API：

cudaGraphicsGLRegisterImage();
cudaGraphicsSubResourceGetMappedArray();

　　VBO/PBO是屬於OpenGL Buffer對象，而OpenGL Texture則是另一種對象。因此，兩種類型的處理需要區別對待。在這個地方耽擱了很久，就是因為沒有看文檔說明。下面一段話正是對這種情況的說明：

From the CUDA Reference Guide entry for `cudaGraphicsResourceGetMappedPointer()`:

> If resource is not a buffer then it cannot be accessed via a pointer and cudaErrorUnknown is returned.

From the CUDA Reference Guide entry for `cudaGraphicsSubResourceGetMappedArray()`:

> If resource is not a texture then it cannot be accessed via an array and cudaErrorUnknown is returned.

In other words, use **GetMappedPointer** for mapped buffer objects. Use **GetMappedArray** for mapped texture objects.

三、參考鏈接

• http://stackmirror.cn/page/4ejhmgxan1w
• https://stackoverflow.com/questions/21765604/draw-image-from-vertex-buffer-object-generated-with-cuda-using-opengl
• https://stackoverflow.com/questions/19244191/cuda-opengl-interop-draw-to-opengl-texture-with-cuda?rq=1
• https://www.3dgep.com/opengl-interoperability-with-cuda/