Odd Misaligned Error with Surface Object

Accelerated Computing CUDA CUDA Programming and Performance
1

Hello, I’m currently trying to interop Cuda and OpenGL on my project but it’s proving to be a challenging task. I’m currently using the following code:

GLuint _frameTextureGL;
cudaGraphicsResource_t _frameTextureCUDA;

void GLInit(){
  glEnable(GL_TEXTURE_2D);
  glGenTextures(1, &_frameTextureGL);

  glBindTexture(GL_TEXTURE_2D, _frameTextureGL);
  {
    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, RES_X, RES_Y, 0, GL_RGBA, GL_FLOAT, NULL);
  }
  glBindTexture(GL_TEXTURE_2D, 0);

  CudaErrorCheck(cudaGraphicsGLRegisterImage(&_frameTextureCUDA, _frameTextureGL, GL_TEXTURE_2D, cudaGraphicsRegisterFlagsWriteDiscard));

}

void DrawFrame(){

  LaunchKernel();

  glBindTexture(GL_TEXTURE_2D, _frameTextureGL);
  {
    glBegin(GL_QUADS);
    {
      glTexCoord2f(0.0f, 0.0f);
      glVertex2f(0.0f, 0.0f);
      glTexCoord2f(1.0f, 0.0f);
      glVertex2f(1.0f, 0.0f);
      glTexCoord2f(1.0f, 1.0f);
      glVertex2f(1.0f, 1.0f);
      glTexCoord2f(0.0f, 1.0f);
      glVertex2f(0.0f, 1.0f);
    }
    glEnd();
  }
  glBindTexture(GL_TEXTURE_2D, 0);

  glFinish();
  glutPostRedisplay();
}

LaunchKernel(){

  CudaErrorCheck (cudaGraphicsMapResources(1, &_frameTextureCUDA));  //error happens here
  
  cudaArray_t frameCudaArray;
  CudaErrorCheck( cudaGraphicsSubResourceGetMappedArray(&frameCudaArray, _frameTextureCUDA, 0, 0));
  cudaResourceDesc frameCudaArrayResourceDesc;
  
  frameCudaArrayResourceDesc.resType = cudaResourceTypeArray;
  frameCudaArrayResourceDesc.res.array.array = frameCudaArray;
    
  cudaSurfaceObject_t frameCudaSurfaceObject;
  CudaErrorCheck( cudaCreateSurfaceObject(&frameCudaSurfaceObject, &frameCudaArrayResourceDesc));
    
  KernelFunction<< <(RES_X * RES_Y / CUDA_THREADS_PER_BLOCK) + 1, CUDA_THREADS_PER_BLOCK >> >( frameCudaSurfaceObject);
    
  CudaErrorCheck( cudaDestroySurfaceObject(frameCudaSurfaceObject));
  
  CudaErrorCheck( cudaGraphicsUnmapResources(1, &_frameTextureCUDA));
  cudaStreamSynchronize(0);

}

__global__ void KernelFunction(cudaSurfaceObject_t frameCudaSurfaceObject){

  const int maxIndex = RES_X * RES_Y;
  int idx = threadIdx.x + blockDim.x * blockIdx.x;

  if (idx > maxIndex) return;

  surf2Dwrite(make_float4(1.0f, 0.0f, 0.0f, 1.0f), frameCudaSurfaceObject, 0, 0);

  //surf2Dwrite(make_float4(1.0f, 0.0f, 0.0f, 1.0f), frameCudaSurfaceObject, (idx % RES_X) * 4, idx / RES_X);
  
}

The problem is that the surf2Dwrite on the kernel will cause other parts of my code to get an error.

With the code on it’s current state my project runs but it doesn’t show anything (which I suppose is normal since I’m not filling the surface object correctly).

But if I change my X & Y coordinates to anything other than (0,0) (by uncommenting the last line for example) then I get a “missaligned error” on the first line of the LaunchKernel function. (this error only shows on the second frame, not the first)

Any ideia what might be causing this? I’ve looked though countless samples of code and stackoverflow pages and can’t seem to figure out what the problem is with my code.

2

Turns out surd2Dwrite uses coordinates in bytes and not in elements.

So for float4 I should have been using

surf2Dwrite(make_float4(1.0f, 0.0f, 0.0f, 1.0f), frameCudaSurfaceObject, (idx % RES_X) * 16, idx / RES_X);