title: Cuda + OpenGL interoperability , glDrawArrays() Access violation
body: <p>first I'd like to apologize for the amount of code I'm providing below, I'm not sure but I could be infringing the posting rules with that.</p>
  
  <p>I'm trying to modify nvidia's cuda sample fluidsGL (example can be seen here <a href="https://www.youtube.com/watch?v=jU9wgSu4_zw" rel="nofollow">https://www.youtube.com/watch?v=jU9wgSu4_zw</a>) so that instead of modifying random pixels, I could load a custom image such as a PNG, and then apply the fluid solver over it,effectively distorting it in a fluid-like fashion.</p>
  
  <p>After trying lots of different implementations, it ocurred me that it would probably be possible by using two vbos, one for the float2 particles that will be sent to the fluid solver kernels, and one for the colors, which I fill in with the RGBA information from the image.</p>
  
  <p>I'm getting an access violation within glDrawArrays though, and I still havent figured out why, in earlier attempts I tried to use a single vbo with a struct called vertex_data, that held a float2 and an uchar4, sent it via glbufferdata and let the kernel calculate over the float2 portion of it, but the access violation exception happened as well.</p>
  
  <p>If anyone is willing to help me with that, I'd be grateful</p>
  
  <pre><code>typedef unsigned char ubyte;
  
  #define DEVICE __device__
  #define GLOBAL __global__
  
  #define MAX(a,b) ((a &gt; b) ? a : b)
  #define DIM 512
  #define DS DIM*DIM
  
  glm::mat4 m_mat;
  
  // CUFFT plan handle
  cufftHandle planr2c;
  cufftHandle planc2r;
  static float2 *vxfield = NULL;
  static float2 *vyfield = NULL;
  
  float2 *hvfield = NULL;
  float2 *dvfield = NULL;
  static int wWidth = MAX(512, DIM);
  static int wHeight = MAX(512, DIM);
  
  /*-----CUSTOM STRUCT-----------------------------------------------------*/
  
  struct GLTexture
  {
      GLuint id;
      int width;
      int height;
  };
  
  
  vertex_data data[DS];
  
  //c linkage
  /*--------------------------------------------------------------------------------------------------------------------------------*/
  extern "C" void addForces(float2 *v, int dx, int dy, int spx, int spy, float fx, float fy, int r);
  extern "C" void advectVelocity(float2 *v, float *vx, float *vy, int dx, int pdx, int dy, float dt);
  extern "C" void diffuseProject(float2 *vx, float2 *vy, int dx, int dy, float dt, float visc);
  extern "C" void updateVelocity(float2 *v, float *vx, float *vy, int dx, int pdx, int dy);
  extern "C" void advectParticles(GLuint vbo, float2 *v, int dx, int dy, float dt);
  /*--------------------------------------------------------------------------------------------------------------------------------*/
  
  GLSLProgram prog;
  IOManager m_manager;
  GLTexture m_tex;
  std::vector&lt;ubyte&gt; in_img;
  std::vector&lt;ubyte&gt; out_img;
  vertex_data vData[6];
  
  
  GLuint positionsVBO;
  GLuint colorsVBO;
  cudaGraphicsResource* positionsVBO_CUDA;
  float2 *particles = NULL;
  
  float2 *part_cuda = NULL;
  
  int lastx = 0, lasty = 0;
  int clicked = 0;
  size_t tPitch = 0;
  
  float myrand(void)
  {
      return rand() / (float)RAND_MAX;
  }
  
  void initParticles(float2 *p, int dx, int dy)
  {
      int i, j;
  
      for (i = 0; i &lt; dy; i++)
      {
          for (j = 0; j &lt; dx; j++)
          {
              p[i*dx + j].x = (j + 0.5f + (myrand() - 0.5f)) / dx;
              p[i*dx + j].y = (i + 0.5f + (myrand() - 0.5f)) / dy;
          }
      }
  
  }
  
  
  
  
  void keyboard(unsigned char key, int x, int y)
  {
      switch (key)
      {
      case 27:
  
          glutDestroyWindow(glutGetWindow());
          exit(0);
          return;
  
  
      default:
          break;
      }
  }
  
  void click(int button, int updown, int x, int y)
  {
      lastx = x;
      lasty = y;
      clicked = !clicked;
  }
  
  void motion(int x, int y)
  {
      // Convert motion coordinates to domain
      float fx = (lastx / (float)wWidth);
      float fy = (lasty / (float)wHeight);
      int nx = (int)(fx * DIM);
      int ny = (int)(fy * DIM);
  
      if (clicked &amp;&amp; nx &lt; DIM - FR &amp;&amp; nx &gt; FR - 1 &amp;&amp; ny &lt; DIM - FR &amp;&amp; ny &gt; FR - 1)
      {
          int ddx = x - lastx;
          int ddy = y - lasty;
          fx = ddx / (float)wWidth;
          fy = ddy / (float)wHeight;
          int spy = ny - FR;
          int spx = nx - FR;
          addForces(dvfield, DIM, DIM, spx, spy, FORCE * DT * fx, FORCE * DT * fy, FR);
          lastx = x;
          lasty = y;
      }
  
      glutPostRedisplay();
  }
  
  void cleanup(void)
  {
      cudaGraphicsUnregisterResource(positionsVBO_CUDA);
  
      unbindTexture();
      deleteTexture();
  
      // Free all host and device resources
      free(hvfield);
      free(particles);
      cudaFree(dvfield);
      cudaFree(vxfield);
      cudaFree(vyfield);
      cufftDestroy(planr2c);
      cufftDestroy(planc2r);
  
      glBindBufferARB(GL_ARRAY_BUFFER_ARB, 0);
      glDeleteBuffersARB(1, &amp;positionsVBO);
  }
  
  void run_simulation(void){
  
      //Execute kernel
  
      advectVelocity(dvfield, (float *)vxfield, (float *)vyfield, DIM, RPADW, DIM, DT);
      diffuseProject(vxfield, vyfield, CPADW, DIM, DT, VIS);
      updateVelocity(dvfield, (float *)vxfield, (float *)vyfield, DIM, RPADW, DIM);
      advectParticles(positionsVBO, dvfield, DIM, DIM, DT);
  
  
  
  }
  
  
  void initShaders(){
      prog.compileShaders("vShader.vertex", "fShader.frag");
      prog.addAttribute("vertexPos");
      prog.addAttribute("vertexColor");
  
      prog.linkShaders();
  }
  
  
  
  void pre_display()
  {
  
      glViewport(0, 0, 512, 512);
      glutPostRedisplay();
  
  }
  
  void display()
  {
      pre_display();
  
      // render points from vertex buffer
      glClear(GL_COLOR_BUFFER_BIT);
  
      initShaders();
  
      run_simulation();
  
      prog.use();
  
      //GLint textureUniform = prog.getUniformLocation("mySampler");
      //glUniform1i(textureUniform, 0);
      //glActiveTexture(GL_TEXTURE0);
  
      GLint pUniform = prog.getUniformLocation("P");
      glUniformMatrix4fv(pUniform, 1, GL_FALSE, &amp;m_mat[0][0]);
  
      glBindBufferARB(GL_ARRAY_BUFFER_ARB, positionsVBO);
  
      glPointSize(1);
  
      glEnable(GL_POINT_SMOOTH);
      glEnable(GL_BLEND);
      glBlendFunc(GL_SRC_ALPHA, GL_ONE_MINUS_SRC_ALPHA);
  
      //glEnableVertexAttribArray(0); tried to manually enable the arrays
      //glEnableVertexAttribArray(1);
  
  
      glDisable(GL_DEPTH_TEST);
      glDisable(GL_CULL_FACE);
  
      glVertexAttribPointer(0, 2, GL_FLOAT, GL_TRUE, sizeof(float2), 0);
  
      glBindBufferARB(GL_ARRAY_BUFFER_ARB, 0); //dont need this but access violaton persists without it anyway
  
      glBindBufferARB(GL_ARRAY_BUFFER_ARB, colorsVBO);
  
      glVertexAttribPointer(1, 4, GL_UNSIGNED_BYTE, GL_TRUE, sizeof(vertex_data) ,(void*)(offsetof(vertex_data, col)));
  
  
      glDrawArrays(GL_POINTS, 0, DS);
  
      glBindBufferARB(GL_ARRAY_BUFFER_ARB, 0);
  
      prog.unuse();
  
  
      glDisable(GL_TEXTURE_2D); //from nvidia's probably linked to the cudaarray_t
  
  
      // Swap buffers
      glutSwapBuffers();
  
  }
  
  void initGL()
  {
      int foo = 1;
      char *bar = "bar";
      glutInit(&amp;foo, &amp;bar);
      glutInitDisplayMode(GLUT_DEPTH | GLUT_DOUBLE | GLUT_RGBA);
      glutInitWindowSize(DIM, DIM);
      glutCreateWindow("mate21");
  
      glClearColor(0.0, 0.0, 0.0, 1.0);
  
  
      glutKeyboardFunc(keyboard);
      glutMouseFunc(click);
      glutMotionFunc(motion);
  
  
  
      glutDisplayFunc(display);
  
      glewInit();
  }
  
  void setGLDevice(){
      cudaDeviceProp prop;
      int dev;
  
      memset(&amp;prop, 0, sizeof(cudaDeviceProp));
      prop.major = 1;
      prop.minor = 0;
      cudaChooseDevice(&amp;dev, &amp;prop);
      cudaGLSetGLDevice(dev);
  }
  
  
  
  void createVBO(){
  
      //reading rgba information from image to out_img
      unsigned long width, height;
      IOManager::readFileToBuffer("jojo_test.png", in_img);
      decodePNG(out_img, width, height, &amp;(in_img[0]), in_img.size());
  
  
      //data.resize(DS); data used to be a vector, gave up on that
      for (int i = 0; i &lt; DS; ++i){
  
          //data[i].pos = particles[i]; edited vertex_data struct for rgba only
          data[i].col.x = out_img[i * 4 + 0];
          data[i].col.y = out_img[i * 4 + 1];
          data[i].col.z = out_img[i * 4 + 2];
          data[i].col.w = out_img[i * 4 + 3];
  
      }
  
  
      glGenBuffers(1, &amp;positionsVBO);
      glBindBufferARB(GL_ARRAY_BUFFER_ARB, positionsVBO);
  
      glBufferDataARB(GL_ARRAY_BUFFER_ARB, sizeof(float2)*DS, particles , GL_DYNAMIC_DRAW_ARB);
      glBindBufferARB(GL_ARRAY_BUFFER_ARB, 0);
  
      cudaGraphicsGLRegisterBuffer(&amp;positionsVBO_CUDA, positionsVBO, cudaGraphicsMapFlagsNone);
  
      glBindBufferARB(GL_ARRAY_BUFFER_ARB, 0);
  
      glGenBuffers(1, &amp;colorsVBO);
      glBindBuffer(GL_ARRAY_BUFFER_ARB, colorsVBO);
      glBufferDataARB(GL_ARRAY_BUFFER_ARB, sizeof(vertex_data)*DS, data, GL_DYNAMIC_DRAW_ARB);
      glBindBufferARB(GL_ARRAY_BUFFER_ARB, 0);
  
  
  
  }
  
  int main()
  {
      setGLDevice();
  
      initGL();
  
      //orthogonal view matrix with glm
      m_mat = glm::ortho(0, 1, 1, 0, 0, 1);
  
  
  
  
  
      hvfield = (float2 *)malloc(sizeof(float2) * DS);
      memset(hvfield, 0, sizeof(float2) * DS);
  
      // Allocate and initialize device data
      cudaMallocPitch((void **)&amp;dvfield, &amp;tPitch, sizeof(float2)*DIM, DIM);
  
      cudaMemcpy(dvfield, hvfield, sizeof(float2) * DS,
          cudaMemcpyHostToDevice);
      // Temporary complex velocity field data
      cudaMalloc((void **)&amp;vxfield, sizeof(float2) * PDS);
      cudaMalloc((void **)&amp;vyfield, sizeof(float2) * PDS);
  
      setupTexture(DIM, DIM);
      bindTexture();
  
      // Create particle array
      particles = (float2 *)malloc(sizeof(float2) * DS);
      memset(particles, 0, sizeof(float2) * DS);
  
      initParticles(particles, DIM, DIM);
  
      // Create CUFFT transform plan configuration
      cufftPlan2d(&amp;planr2c, DIM, DIM, CUFFT_R2C);
      cufftPlan2d(&amp;planc2r, DIM, DIM, CUFFT_C2R);
  
      cufftSetCompatibilityMode(planr2c, CUFFT_COMPATIBILITY_FFTW_PADDING);
      cufftSetCompatibilityMode(planc2r, CUFFT_COMPATIBILITY_FFTW_PADDING);
  
  
      createVBO();
  
      //cleanup
      glutCloseFunc(cleanup);
      //Launch rendering loop
      glutMainLoop();
  }
  </code></pre>
  
  <p>And this is the relevant kernel that operates on the float2 particles effectively simulating the fluid</p>
  
  <pre><code>extern "C"
  void advectParticles(GLuint vbo, float2 *v, int dx, int dy, float dt)
  {
      dim3 grid((dx/TILEX)+(!(dx%TILEX)?0:1), (dy/TILEY)+(!(dy%TILEY)?0:1));
      dim3 tids(TIDSX, TIDSY);
  
      float2 *p;
      cudaGraphicsMapResources(1, &amp;positionsVBO_CUDA, 0);
  
  
     size_t num_bytes;
     cudaGraphicsResourceGetMappedPointer((void **)&amp;p, &amp;num_bytes,positionsVBO_CUDA);
  
  
      advectParticles_k&lt;&lt;&lt;grid, tids&gt;&gt;&gt;(p, v, dx, dy, dt, TILEY/TIDSY, tPitch);
  
  
      cudaGraphicsUnmapResources(1, &amp;positionsVBO_CUDA, 0);
  
  }
  </code></pre>
  
  <p>And advectParticles_k:</p>
  
  <pre><code>__global__ void
  advectParticles_k(float2 *part, float2 *v, int dx, int dy,
                    float dt, int lb, size_t pitch)
  {
  
      int gtidx = blockIdx.x * blockDim.x + threadIdx.x;
      int gtidy = blockIdx.y * (lb * blockDim.y) + threadIdx.y * lb;
      int p;
  
      // gtidx is the domain location in x for this thread
      float2 pterm, vterm;
  
      if (gtidx &lt; dx)
      {
          for (p = 0; p &lt; lb; p++)
          {
              // fi is the domain location in y for this thread
              int fi = gtidy + p;
  
              if (fi &lt; dy)
              {
                  int fj = fi * dx + gtidx;
                  pterm = part[fj];
  
                  int xvi = ((int)(pterm.x * dx));
                  int yvi = ((int)(pterm.y * dy));
                  vterm = *((float2 *)((char *)v + yvi * pitch) + xvi);
  
                  pterm.x += dt * vterm.x;
                  pterm.x = pterm.x - (int)pterm.x;
                  pterm.x += 1.f;
                  pterm.x = pterm.x - (int)pterm.x;
                  pterm.y += dt * vterm.y;
                  pterm.y = pterm.y - (int)pterm.y;
                  pterm.y += 1.f;
                  pterm.y = pterm.y - (int)pterm.y;
  
                  part[fj] = pterm;
              }
          } // If this thread is inside the domain in Y
      } // If this thread is inside the domain in X
  }
  </code></pre>
  
  <p>Thank you</p>