OpenGL texture interaction with surface object API without copy

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi!
I actually have some nice part of working code but I would still like to know some comments about it, if this could be achieved easier. My goal is to write data to an OpenGL-Texture which is calculated through heavy CUDA processing. I already managed to get another example working which uses pitched 2D memory and cudaMemcpy2DToArray.

By using surface object API, I want to avoid this memory copy. So here my example (just the relevant parts):

cudaArray *textureCudaArray;
struct cudaResourceDesc resDesc;
memset(&resDesc, 0, sizeof(resDesc));
resDesc.resType = cudaResourceTypeArray;
cudaSurfaceObject_t deviceTextureGraphicSurface = 0;

cudaGraphicsGLRegisterImage(...)

// start render loop
while(...)
{
  ...
  cudaGraphicsMapResources(1, &textureGraphicResource, 0);
  cudaGraphicsSubResourceGetMappedArray(&textureCudaArray, textureGraphicResource, 0, 0);
  resDesc.res.array.array = textureCudaArray;
  cudaCreateSurfaceObject(&deviceTextureGraphicSurface, &resDesc);
  myTextureKernel<<<16, 32>>>(deviceTextureGraphicSurface, textureWidth, textureHeight);
  checkCuda( cudaGraphicsUnmapResources(1, &textureGraphicResource, 0) );
  ...
}
...
cudaGraphicsUnregisterResource(textureGraphicResource);
cudaDestroySurfaceObject(deviceTextureGraphicSurface);

Now some questions about it:

  1. Is data really copied during cudaCreateSurfaceObject? Or is a pointer mapped to the data of the cudaArray?
  2. While I looked for the solution I came across the C++ API for this https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__HIGHLEVEL.html#group__CUDART__HIGHLEVEL with a function cudaBindSurfaceToArray. Do these routines just map to the deprecated C routines?
  3. Is there a way to do this task better? Maybe with Texture Object API? But I read, that textures are just used to read data, not to constantly write data? Am I correct?