I am trying to do some image processing on opengl texture and then use the output texture into opengl shader.
The problem I am facing is , I am getting blank texture after the surface write in cuda. Strange thing is , when I display the processed texture right after kernel call , by copying texture to host memory using cudaMemcpyFromArray and I get correct results. But when I try to get same texture using glGetTexImage , I get blank result. I have verified further that the written values are fine , because when I access the same texture in other cuda kernel , the values are correct. It is just glGetTexImage , which is behaving incorrectly and it’s not reporting any opengl error either. What could be possible reason for this? :
Here is the code for mapping the texture to cuda ( and calling the cuda kernel ):
registerGLResources();
HANDLE_ERROR( cudaGraphicsMapResources( 2 , mResources + 6 ) );
cudaArray_t imarray , imDerArray ;
HANDLE_ERROR( cudaGraphicsSubResourceGetMappedArray( &imarray , mResources[ 6 ] , 0 , 0 ) );
HANDLE_ERROR( cudaGraphicsSubResourceGetMappedArray( &imDerArray , mResources[ 7 ] , 0 , 0 ) );
// Specify texture struct
cudaResourceDesc resDesc1 , resDesc2 ;
memset(&resDesc1, 0, sizeof(resDesc1));
resDesc1.resType = cudaResourceTypeArray;
resDesc1.res.array.array = imarray;
memset(&resDesc2, 0, sizeof(resDesc2));
resDesc2.resType = cudaResourceTypeArray;
resDesc2.res.array.array = imDerArray;
cudaSurfaceObject_t imder = 0;
cudaTextureObject_t im = 0;
// Specify texture object parameters struct
cudaTextureDesc texDesc1;
memset(&texDesc1, 0, sizeof(texDesc1));
texDesc1.addressMode[0] = cudaAddressModeWrap;
texDesc1.addressMode[1] = cudaAddressModeWrap;
texDesc1.filterMode = cudaFilterModeLinear;
texDesc1.readMode = cudaReadModeElementType;
texDesc1.normalizedCoords = 0;
HANDLE_ERROR( cudaCreateTextureObject( &im , &resDesc1, &texDesc1, NULL) );
HANDLE_ERROR( cudaCreateSurfaceObject( &imder, &resDesc2 ) );
std::cout << " computing image derivatives " << mNeighborWidth << " " << mNeighborHeight << std::endl;
HANDLE_ERROR( vc::mvg::computeImageDerivatives( im , imder , mNeighborWidth , mNeighborHeight ) );
cv::Mat temp1(mNeighborHeight, mNeighborWidth, CV_32FC2), temp2(mNeighborHeight, mNeighborWidth, CV_32FC1);
HANDLE_ERROR(cudaMemcpyFromArray(temp1.data, imDerArray, 0, 0, 2 * mNeighborWidth * mNeighborHeight * sizeof(float), cudaMemcpyDeviceToHost));
Eigen::Vector2f *t1 = (Eigen::Vector2f*)temp1.data;
float *t2 = (float*)temp2.data;
for (int yy = 0; yy < mNeighborHeight; yy++)
for (int xx = 0; xx < mNeighborWidth; xx++)
{
t2[yy * mNeighborWidth + xx] = t1[yy * mNeighborWidth + xx](0);
}
cv::namedWindow("derivative1", 0);
cv::imshow("derivative1", temp2);
cv::waitKey();
HANDLE_ERROR( cudaDestroyTextureObject( im ) );
HANDLE_ERROR( cudaDestroySurfaceObject( imder ) );
HANDLE_ERROR( cudaGraphicsUnmapResources( 2 , mResources + 6 , NULL ) );
cv::Mat temp3(mNeighborHeight, mNeighborWidth, CV_32FC2), temp4(mNeighborHeight, mNeighborWidth, CV_32FC1);
// //
GL_CHECK(glBindTexture(GL_TEXTURE_2D, mImageDerivative));
GL_CHECK(glGetTexImage(GL_TEXTURE_2D, 0, GL_RG, GL_FLOAT, temp3.data));
GL_CHECK(glBindTexture(GL_TEXTURE_2D, 0));
Eigen::Vector2f *t3 = (Eigen::Vector2f*)temp3.data;
float *t4 = (float*)temp4.data;
for (int yy = 0; yy < mNeighborHeight; yy++)
for (int xx = 0; xx < mNeighborWidth; xx++)
{
t4[yy * mNeighborWidth + xx] = t3[yy * mNeighborWidth + xx](0);
}
cv::namedWindow("derivative2", 0);
cv::imshow("derivative2", temp4);
cv::waitKey();
I register the the opengl resources correctly. Here is the code for registering the texture for cuda inter operability:
void MVVSurfaceRefinerOGL33CUDABuffer::registerGLResources()
{
if (!mResourcesRegistered)
{
HANDLE_ERROR(cudaGraphicsGLRegisterImage(mResources, mRenderedDepths, GL_TEXTURE_2D, cudaGraphicsMapFlagsNone));
HANDLE_ERROR(cudaGraphicsGLRegisterImage((mResources + 1), mRenderedMask, GL_TEXTURE_2D, cudaGraphicsMapFlagsNone));
HANDLE_ERROR(cudaGraphicsGLRegisterImage((mResources + 2), mColorReferenceTexture, GL_TEXTURE_2D, cudaGraphicsMapFlagsNone));
HANDLE_ERROR(cudaGraphicsGLRegisterImage((mResources + 3), mRenderedTexture, GL_TEXTURE_2D, cudaGraphicsMapFlagsNone));
HANDLE_ERROR(cudaGraphicsGLRegisterImage((mResources + 4), mCost, GL_TEXTURE_2D, cudaGraphicsMapFlagsNone));
HANDLE_ERROR(cudaGraphicsGLRegisterImage((mResources + 5), mCostDer, GL_TEXTURE_2D, cudaGraphicsMapFlagsNone));
HANDLE_ERROR(cudaGraphicsGLRegisterImage((mResources + 6), mTexture, GL_TEXTURE_2D, cudaGraphicsMapFlagsNone));
HANDLE_ERROR(cudaGraphicsGLRegisterImage((mResources + 7), mImageDerivative, GL_TEXTURE_2D, cudaGraphicsMapFlagsNone));
HANDLE_ERROR(cudaGraphicsGLRegisterImage((mResources + 8), mBaryCentricCoords, GL_TEXTURE_2D, cudaGraphicsMapFlagsNone));
HANDLE_ERROR(cudaGraphicsGLRegisterImage((mResources + 9), mDerCoeffs, GL_TEXTURE_2D, cudaGraphicsMapFlagsNone));
HANDLE_ERROR(cudaGraphicsGLRegisterImage((mResources + 10), mFaceIds, GL_TEXTURE_2D, cudaGraphicsMapFlagsNone));
HANDLE_ERROR(cudaGraphicsGLRegisterBuffer((mResources + 11), mPhotometricDerVBO, cudaGraphicsMapFlagsNone));
}
mResourcesRegistered = true;
// std::cout<<" opengl resourced registered "<<std::endl;
}
Here is the code for texture creation in opengl:
GL_CHECK(glBindTexture(GL_TEXTURE_2D, mImageDerivative));
GL_CHECK(glTexImage2D(GL_TEXTURE_2D, 0, GL_RG32F, mNeighborWidth,
mNeighborHeight, 0, GL_RG, GL_FLOAT, 0));
}
GL_CHECK(glBindTexture(GL_TEXTURE_2D, 0));
Following is the link for outputs of code:
http://i.stack.imgur.com/F2qBv.png
http://i.stack.imgur.com/myU8j.png