Cuda + OpenGL互操作性，glDrawArrays（）訪問衝突

首先，我想爲我在下面提供的代碼數量表示歉意，我不確定，但是我可能會因此違反發佈規則。Cuda + OpenGL互操作性，glDrawArrays（）訪問衝突

我試圖修改nvidia的cuda示例流體GL（例如可以在這裏看到https://www.youtube.com/watch?v=jU9wgSu4_zw），這樣就不需要修改隨機像素，我可以加載一個自定義圖像，比如PNG，然後在其上應用流體求解器，以流體般的方式有效地扭曲它。

嘗試了很多不同的實現之後，它使我可能通過使用兩個vbos，一個用於float2粒子將被髮送到流體求解器內核，另一個用於我填充的顏色與來自圖像的RGBA信息。

雖然我在glDrawArrays中發現了一個訪問衝突，但我仍然沒弄清楚爲什麼，在之前的嘗試中，我嘗試使用一個名爲vertex_data的結構體，它包含一個float2和一個uchar4， glbufferdata並讓內核計算它的float2部分，但是訪問衝突異常也發生了。

如果有人願意幫助我，我會很感激

typedef unsigned char ubyte; 

#define DEVICE __device__ 
#define GLOBAL __global__ 

#define MAX(a,b) ((a > 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<ubyte> in_img; 
std::vector<ubyte> 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 < dy; i++) 
    { 
     for (j = 0; j < 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 && nx < DIM - FR && nx > FR - 1 && ny < DIM - FR && ny > 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, &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, &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(&foo, &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(&prop, 0, sizeof(cudaDeviceProp)); 
    prop.major = 1; 
    prop.minor = 0; 
    cudaChooseDevice(&dev, &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, &(in_img[0]), in_img.size()); 


    //data.resize(DS); data used to be a vector, gave up on that 
    for (int i = 0; i < 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, &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(&positionsVBO_CUDA, positionsVBO, cudaGraphicsMapFlagsNone); 

    glBindBufferARB(GL_ARRAY_BUFFER_ARB, 0); 

    glGenBuffers(1, &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 **)&dvfield, &tPitch, sizeof(float2)*DIM, DIM); 

    cudaMemcpy(dvfield, hvfield, sizeof(float2) * DS, 
     cudaMemcpyHostToDevice); 
    // Temporary complex velocity field data 
    cudaMalloc((void **)&vxfield, sizeof(float2) * PDS); 
    cudaMalloc((void **)&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(&planr2c, DIM, DIM, CUFFT_R2C); 
    cufftPlan2d(&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(); 
}

而且這是在FLOAT2顆粒有效運行模擬流體

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, &positionsVBO_CUDA, 0); 


    size_t num_bytes; 
    cudaGraphicsResourceGetMappedPointer((void **)&p, &num_bytes,positionsVBO_CUDA); 


    advectParticles_k<<<grid, tids>>>(p, v, dx, dy, dt, TILEY/TIDSY, tPitch); 


    cudaGraphicsUnmapResources(1, &positionsVBO_CUDA, 0); 

}

而且相關的內核advectParticles_k：

__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 < dx) 
    { 
     for (p = 0; p < lb; p++) 
     { 
      // fi is the domain location in y for this thread 
      int fi = gtidy + p; 

      if (fi < 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 
}

謝謝

來源

2016-07-29 Kareshi

我覺得你沒有在緩衝區中爲你的粒子分配足夠的內存。

讓我們假設：

sizeof(float2) = 4 
DS = 10

調用此：

glBufferDataARB(GL_ARRAY_BUFFER_ARB, sizeof(float2)*DS, particles , GL_DYNAMIC_DRAW_ARB);

您分配了4 * 10 = 字節的內存

然後你調用這個函數：

glVertexAttribPointer(0, 2, GL_FLOAT, GL_TRUE, sizeof(float2), 0);

其中說，你需要2 * 4 = 單個粒子的內存字節。

然後要渲染：

glDrawArrays(GL_POINTS, 0, DS);

有了這個叫你正在努力使10個顆粒，其消耗10 * 8 = 字節的內存

而與此叫你得到訪問衝突，因爲您只分配了字節的內存。

來源

2016-07-29 20:30:09 Amadeusz

嗯，我看到，雖然第一眼看到，但float2 cuda本地類型是結構，它包含兩個浮點值x和y。對於opengl的兼容性，它應該可以在glVertexAttribPointer調用中正常工作，我會考慮你的評論並嘗試進一步檢查它。如果我設法取得任何積極成果，我會再次發帖。謝謝Amadeusz。 – Kareshi

爲什麼你會認爲'sizeof（float2）= 4'？它是8，並且一個微不足道的測試將證明它。 –

Cuda + OpenGL互操作性，glDrawArrays（）訪問衝突

回答

相關問題