Bug with CUDA / OpenGL Interop (see https://stackoverflow.com/questions/53901132/cuda-opengl-interop-writing-to-surface-object-does-not-erase-previous-contents)

fragment.fs

#version 330 core
out vec4 FragColor;
in vec2 textureCoord;

uniform sampler2D textureData;

void main(){
	FragColor = texture(textureData, textureCoord);
}

main.cu

// graphics
#include <glad/glad.h>
#include <GLFW/glfw3.h>
// standard library
#include <stdlib.h>
#include <stdio.h>
#include <iostream>
#include <fstream>
#include <string>
#include <math.h>
// cuda
#include <cuda_gl_interop.h>
#include <cuda.h>
#include <cuda_runtime.h>
// project files
#include "shader.h"
#include "stb_image.h"

using namespace std;

// -----------------------------------------------------------------------------

// window dimensions
const int WIDTH = 800;
const int HEIGHT = 800;

// image properties
int g_imageWidth;
int g_imageHeight;

// graphics interop
struct cudaResourceDesc g_resourceDesc;
cudaGraphicsResource_t g_textureResource;
cudaArray_t g_textureArray;
cudaSurfaceObject_t g_surfaceObj = 0;

// HANDLE ERRORS ---------------------------------------------------------------

static void HandleError( cudaError_t err, const char *file, int line ) {
    if (err != cudaSuccess) {
        printf( "%s in %s at line %d\n", cudaGetErrorString( err ), file, line );
        exit( EXIT_FAILURE );
    }
}

#define HANDLE( err ) (HandleError( err, __FILE__, __LINE__ ))

#define HANDLE_NULL( a ) {if (a == NULL) { \
                            printf( "Host memory failed in %s at line %d\n", \
                                    __FILE__, __LINE__ ); \
                            exit( EXIT_FAILURE );}}

// GEOMETRY --------------------------------------------------------------------

float vertices[] = {
    // positions          // colors           // texture coords
     0.5f,  0.5f, 0.0f,   1.0f, 0.0f, 0.0f,   1.0f, 1.0f,   // top right
     0.5f, -0.5f, 0.0f,   0.0f, 1.0f, 0.0f,   1.0f, 0.0f,   // bottom right
    -0.5f, -0.5f, 0.0f,   0.0f, 0.0f, 1.0f,   0.0f, 0.0f,   // bottom left
    -0.5f,  0.5f, 0.0f,   1.0f, 1.0f, 0.0f,   0.0f, 1.0f    // top left 
};

unsigned int indices[] = {  // note that we start from 0!
    0, 1, 3,   // first triangle
    1, 2, 3    // second triangle
};

// OPENGL ----------------------------------------------------------------------

void framebufferSizeCallback(GLFWwindow *window, int width, int height){
    glViewport(0, 0, width, height);
}

GLFWwindow* initWindow(){
    // glfw
    glfwInit();
    glfwWindowHint(GLFW_CONTEXT_VERSION_MAJOR, 3);
    glfwWindowHint(GLFW_CONTEXT_VERSION_MINOR, 3);
    glfwWindowHint(GLFW_OPENGL_PROFILE, GLFW_OPENGL_CORE_PROFILE);

    // create window
    GLFWwindow *window = glfwCreateWindow(WIDTH, HEIGHT, "CUDA/GL Interop", NULL, NULL);
    if(window == NULL) {
        cout << "Failed to create GLFW window!" << endl;
        glfwTerminate();
        std::exit(-1);
    }

    glfwMakeContextCurrent(window);

    // initialize GLAD
    if (!gladLoadGLLoader((GLADloadproc)glfwGetProcAddress)) {
        cout << "Failed to initialize GLAD!" << endl;
        std::exit(-1);
    }

    // clear
    glClearColor(0.0f, 0.0f, 0.0f, 1.0f);
    glDisable(GL_DEPTH_TEST);

    // viewport
    glViewport(0, 0, 800, 600);
    glfwSetFramebufferSizeCallback(window, framebufferSizeCallback);

    return window;
}

void initBuffers(uint *vao, uint *vbo, uint *ebo){
    // vertex array object
    glGenVertexArrays(1, vao);
    glBindVertexArray(*vao);

    // vertex buffer object
    glGenBuffers(1, vbo);
    glBindBuffer(GL_ARRAY_BUFFER, *vbo);
    glBufferData(GL_ARRAY_BUFFER, sizeof(vertices), vertices, GL_STATIC_DRAW);

    // element buffer object
    glGenBuffers(1, ebo);
    glBindBuffer(GL_ELEMENT_ARRAY_BUFFER, *ebo);
    glBufferData(GL_ELEMENT_ARRAY_BUFFER, sizeof(indices), indices, GL_STATIC_DRAW);

    // vertex attribute:  position
    glVertexAttribPointer(0, 3, GL_FLOAT, GL_FALSE, 8*sizeof(float), (void*)0);
    glEnableVertexAttribArray(0);

    // vertex attribute:  texture coords
    glVertexAttribPointer(1, 2, GL_FLOAT, GL_FALSE, 8 * sizeof(float), (void*)(6*sizeof(float)));
    glEnableVertexAttribArray(1);
    
    // unbind (not strictly necessary)
    glBindBuffer(GL_ARRAY_BUFFER, 0);
    glBindVertexArray(0);
}

GLuint initTexturesGL(){
    // load texture from file
    int numChannels;
    unsigned char *data = stbi_load("img/container.jpg", &g_imageWidth, &g_imageHeight, &numChannels, 4);
    if(!data){
        std::cerr << "Error:  Failed to load texture image!" << std::endl;
        exit(1);
    }

    // opengl texture
    GLuint textureId;
    glGenTextures(1, &textureId);
    glBindTexture(GL_TEXTURE_2D, textureId);

    // wrapping
    glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_S, GL_MIRRORED_REPEAT);
    glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_T, GL_MIRRORED_REPEAT);

    // filtering
    glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_LINEAR_MIPMAP_LINEAR);
    glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_LINEAR);

    // set texture image
    glTexImage2D(
        GL_TEXTURE_2D,    // target
        0,                // mipmap level
        GL_RGBA8,         // internal format (#channels, #bits/channel, ...)
        g_imageWidth,     // width
        g_imageHeight,    // height
        0,                // border (must be zero)
        GL_RGBA,          // format of input image
        GL_UNSIGNED_BYTE, // type
        data              // data
    );
    glGenerateMipmap(GL_TEXTURE_2D);

    // unbind and free image
    glBindTexture(GL_TEXTURE_2D, 0);
    stbi_image_free(data);

    return textureId;
}

// CUDA ------------------------------------------------------------------------

int initCudaDevice(){
    // choose cuda device
    cudaDeviceProp deviceProp;
    int deviceId;

    memset(&deviceProp, 0, sizeof(cudaDeviceProp));
    deviceProp.major = 3;
    deviceProp.minor = 0;
    HANDLE( cudaChooseDevice(&deviceId, &deviceProp) );

    return deviceId;
}

void initTexturesCuda(GLuint textureId){
    // register texture
    HANDLE(cudaGraphicsGLRegisterImage(
        &g_textureResource,                       // resource
        textureId,                                // image
        GL_TEXTURE_2D,                            // target
        cudaGraphicsRegisterFlagsSurfaceLoadStore // flags
    ));

    // resource description for surface
    memset(&g_resourceDesc, 0, sizeof(g_resourceDesc));
    g_resourceDesc.resType = cudaResourceTypeArray;
}

__global__ void kernel(cudaSurfaceObject_t surface, int nx, int ny){
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    if(x < nx && y < ny){
        uchar4 data = make_uchar4(x % 255, 
                                  y % 255, 
                                  0, 255);
        surf2Dwrite(data, surface, x * sizeof(uchar4), y);
    }
}

// -----------------------------------------------------------------------------

int main() {
    // CUDA: configure device
    int deviceId = initCudaDevice();
    // GLFW: create window 
    GLFWwindow *window = initWindow();
    // OpenGL: init shader program
    ShaderProgram shader("shaders/vertex.vs", "shaders/fragment.fs");
    // OpenGL: init geometry
    unsigned int vao, vbo, ebo;
    initBuffers(&vao, &vbo, &ebo);
    // OpenGL: init textures
    GLuint textureId = initTexturesGL();
    // CUDA: set up interop
    initTexturesCuda(textureId);

    // thread layout
    dim3 blockDim(32, 32);
    dim3 gridDim((g_imageWidth  + blockDim.x-1)/blockDim.x,
                 (g_imageHeight + blockDim.y-1)/blockDim.y);

    printf("blockDim: %d x %d = %d\n", blockDim.x, blockDim.y, blockDim.x*blockDim.y);
    printf("gridDim: %d x %d = %d\n", gridDim.x, gridDim.y, gridDim.x*gridDim.y);

    // render loop
    while(!glfwWindowShouldClose(window)){
        // -- CUDA --

        // map
        HANDLE(cudaGraphicsMapResources(1, &g_textureResource));


        HANDLE(cudaGraphicsSubResourceGetMappedArray(
            &g_textureArray,   // array through which to access subresource
            g_textureResource, // mapped resource to access
            0,                 // array index
            0                  // mipLevel
        ));

        // create surface object (compute >= 3.0)
        g_resourceDesc.res.array.array = g_textureArray;
        HANDLE(cudaCreateSurfaceObject(&g_surfaceObj, &g_resourceDesc));

        // run kernel
        kernel<<<gridDim, blockDim>>>(g_surfaceObj, g_imageWidth, g_imageHeight);
        
        // unmap
        HANDLE(cudaGraphicsUnmapResources(1, &g_textureResource));

        // --- OpenGL ---

        // clear
        glClear(GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT);

        // use program
        shader.use();

        // triangle
        glBindVertexArray(vao);
        glBindTexture(GL_TEXTURE_2D, textureId);
        glDrawElements(GL_TRIANGLES, 6, GL_UNSIGNED_INT, 0);
        glBindVertexArray(0);

        // glfw:  swap buffers and poll i/o events
        glfwSwapBuffers(window);
        glfwPollEvents();
    }

    // clean up
    glDeleteVertexArrays(1, &vao);
    glDeleteBuffers(1, &vbo);

    // finish
    glfwTerminate();
    return 0;
}

Makefile

# folders
OBJ_DIR := build
OUT_DIR := bin
INC_DIR := include
LIB_DIR := lib
DIRS    := $(OBJ_DIR) $(OUT_DIR) $(INC_DIR) $(LIB_DIR)

# compiler
CC := nvcc
FLAGS := -Wno-deprecated-gpu-targets

OBJECTS := $(addprefix $(OBJ_DIR)/,main.o glad.o stb_image.o)
GL := -lGL -lGLU -lglfw3 -lX11 -lXxf86vm -lXrandr -lpthread -lXi -ldl -lXinerama -lXcursor

.PHONY: all
all: bin/main


$(OBJ_DIR)/main.o: main.cu
	$(CC) -c main.cu -o $@ -I$(INC_DIR) $(FLAGS)

$(OBJ_DIR)/glad.o: $(LIB_DIR)/glad/glad.c
	$(CC) -c $< -o $@ -I$(INC_DIR) $(FLAGS)

$(OBJ_DIR)/stb_image.o: stb_image.c
	$(CC) -c $< -o $@ -I$(INC_DIR) $(FLAGS)

$(OUT_DIR)/main: $(OBJECTS) | $(OUT_DIR)
	$(CC) $(OBJECTS) -o $@ $(GL) $(FLAGS)

# add directories as prerequisites
$(OBJECTS): | $(OBJ_DIR)

# build directory
$(DIRS):
	mkdir $@

# clean
.PHONY: clean
clean:
	rm -rf $(OBJ_DIR)
	rm -rf $(OUT_DIR)

shader.h

#ifndef SHADER_H
#define SHADER_H

#include <glad/glad.h>
#include <GLFW/glfw3.h>
#include <string>
#include <fstream>
#include <sstream>
#include <iostream>

static std::string readFile(const char *filename){
    std::ifstream in(filename);
    std::string contents((std::istreambuf_iterator<char>(in)), std::istreambuf_iterator<char>());
    return contents;
}

static int initVertexShader(){
    // load shaders from file
    std::string shaderString = readFile("shaders/vertex.vs");
    const char* shaderSource = shaderString.c_str();

    // vertex shader
    unsigned int vertexShader;
    vertexShader = glCreateShader(GL_VERTEX_SHADER);
    glShaderSource(vertexShader, 1, &shaderSource, NULL);
    glCompileShader(vertexShader);

    // check errors
    int success;
    char infoLog[512];
    glGetShaderiv(vertexShader, GL_COMPILE_STATUS, &success);

    if(!success) {
        glGetShaderInfoLog(vertexShader, 512, NULL, infoLog);
        std::cerr << "Error:  Failed to compile vertex shader!\n" 
             << infoLog << std::endl;
        exit(-1);
    }

    return vertexShader;
}

static int initFragmentShader(){
    // load shader from file
    const char* shaderSource = readFile("shaders/fragment.fs").c_str();

    // vertex shader
    unsigned int fragShader;
    fragShader = glCreateShader(GL_FRAGMENT_SHADER);
    glShaderSource(fragShader, 1, &shaderSource, NULL);
    glCompileShader(fragShader);

    // check errors
    int success;
    char infoLog[512];
    glGetShaderiv(fragShader, GL_COMPILE_STATUS, &success);

    if(!success) {
        glGetShaderInfoLog(fragShader, 512, NULL, infoLog);
        std::cerr << "Error:  Failed to compile fragment shader!\n" 
             << infoLog << std::endl;
        exit(-1);
    }

    return fragShader;
}


class ShaderProgram {
public:
	GLuint programId;

	// constructor
	ShaderProgram(const GLchar* vertexPath, const GLchar* fragmentPath);

	// use shader
	void use() {
		glUseProgram(this->programId);
	}

	// set uniforms
	void setUniformBool(const std::string &name, bool value) const {
		glUniform1i(glGetUniformLocation(this->programId, name.c_str()), (int)value);
	}
	void setUniformInt(const std::string &name, int value) const {
		glUniform1i(glGetUniformLocation(this->programId, name.c_str()), value);
	}
	void setUniformFloat(const std::string &name, float value) const {
		glUniform1f(glGetUniformLocation(this->programId, name.c_str()), value); 
	}

private:
	static int compileShader(const GLchar *path, const GLenum shaderType);
};

////////////////////////////////////////////////////////////////////////////////

/* SHADERPROGRAM (Constructor)
 * Compiles a vertex and fragment shader, linking them into a ShaderProgram.
 */
ShaderProgram::ShaderProgram(const GLchar* vertexPath, const GLchar* fragmentPath) {
	// compile shaders
	int vertexId = ShaderProgram::compileShader(vertexPath, GL_VERTEX_SHADER);
	int fragmentId = ShaderProgram::compileShader(fragmentPath, GL_FRAGMENT_SHADER);

    //int vertexId = initVertexShader();
    //int fragmentId = initFragmentShader();

	// create and link shader program
	this->programId = glCreateProgram();
    glAttachShader(this->programId, vertexId);
    glAttachShader(this->programId, fragmentId);
    glLinkProgram(this->programId);

    // error check
    int success;
    char infoLog[512];

    glGetProgramiv(this->programId, GL_LINK_STATUS, &success);
    if(!success) {
        glGetProgramInfoLog(this->programId, 512, NULL, infoLog);
        std::cerr << "Error:  Failed to link shader program!\n" 
                  << infoLog << std::endl;
        exit(-1);
    }

    // delete shaders
    glDeleteShader(vertexId);
    glDeleteShader(fragmentId);
}

/* SHADER PROGRAM :: COMPILESHADER
 * Reads shader source from file and compiles a shader
 * of the desired type, with error checking. 
 * @param (path) File location of shader source.
 * @returns (shaderId) OpenGL identifier for the compiled shader.
 */
int ShaderProgram::compileShader(
	const GLchar *path, 
	const GLenum shaderType) 
{
    // create shader
    GLuint shaderId = glCreateShader(shaderType);

	// load and compile shader code
    const char* shaderSource = readFile(path).c_str();
    glShaderSource(shaderId, 1, &shaderSource, NULL);
    glCompileShader(shaderId);

    // check errors
    int success;
    char infoLog[512];
    glGetShaderiv(shaderId, GL_COMPILE_STATUS, &success);

    if(!success) {
        glGetShaderInfoLog(shaderId, 512, NULL, infoLog);
        std::cerr << "Error:  Failed to compile shader!\n" 
                  << infoLog << std::endl;
        exit(-1);
    }

    return shaderId;
}

#endif

stb_image.c

#define STB_IMAGE_IMPLEMENTATION
#include "stb_image.h"

vertex.vs

#version 330 core

layout (location = 0) in vec3 aPos;
layout (location = 1) in vec2 aTex;

out vec2 textureCoord;

void main(){
	gl_Position = vec4(aPos.x, aPos.y, aPos.z, 1.0);
	textureCoord = aTex;
}