CUDA Sampleでよく行われている処理。
ただ、CUDA Sampleはごちゃごちゃしていてよくわからないので、最小構成を抜き出してみました。
int main(int argc, char *argv[])
{
// 画像ファイル入力
cvImg = cv::imread("Lenna.bmp", 0);

// 初期化
init(argc, argv);

glutMainLoop();

// 終了処理
finalize();

return 0;
}
main関数。
画像ファイルの入力と初期化を行った後に、glutMainLoop()を実行します。

void init(int argc, char *argv[])
{
size_t dataSize = cvImg.cols * cvImg.rows * sizeof(unsigned char);

glInit(argc, argv); // GLの初期化
GenTexture(); // テクスチャの生成
GenAndRegisterPBO(dataSize); // PBOの生成とCUDAへの登録
   
// 入力画像のGPU側メモリ確保
  checkCudaErrors(cudaMalloc((unsigned char**)&dImgIn, dataSize));
}
初期化関数です。
以下に各関数を載せます。

void glInit(int argc, char *argv[])
{
glutInit(&argc, argv);
glutInitWindowSize(cvImg.cols, cvImg.rows);
glutInitDisplayMode(GLUT_RGBA | GLUT_DOUBLE);
glutCreateWindow(argv[0]);
glewInit();
glutDisplayFunc(display);
}

void GenTexture()
{
glGenTextures(1, &tex);
glBindTexture(GL_TEXTURE_2D, tex);
glTexImage2D(GL_TEXTURE_2D, 0, GL_LUMINANCE, cvImg.cols, cvImg.rows,
      0, GL_LUMINANCE, GL_UNSIGNED_BYTE, NULL);
glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_NEAREST);
glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST);
}

void GenAndRegisterPBO(const size_t dataSize)
{
glGenBuffers(1, &pbo);
glBindBuffer(GL_PIXEL_UNPACK_BUFFER, pbo);
glBufferData(GL_PIXEL_UNPACK_BUFFER, dataSize, 0, GL_STREAM_DRAW);
glBindBuffer(GL_PIXEL_UNPACK_BUFFER, 0);
cudaGraphicsGLRegisterBuffer(&cuda_pbo_resource, pbo,
      cudaGraphicsMapFlagsWriteDiscard);
}
glInit()はGLUT, GLEWのおなじみの処理。
GenTexture()はテクスチャ生成のおなじみの処理。
GenAndRegisterPBO()はPBOを生成し、そのPBOをCUDAに登録しています。
CUDAへの登録はcudaGraphicsGLRegisterBuffer()で行います。
cudaGraphicsMapFlagsWriteDiscardはPBOがCUDAからWriteOnlyであるというフラグです。

void display()
{
// カーネル関数で書き換える画像(PBO)のポインタを取得
unsigned char *dImgOut;
size_t dataSizeTmp;
checkCudaErrors(cudaGraphicsMapResources(1, &cuda_pbo_resource, 0));
checkCudaErrors(cudaGraphicsResourceGetMappedPointer(
(void **)&dImgOut, &dataSizeTmp, cuda_pbo_resource));

// 入力画像をGPUのグローバルメモリに転送
size_t dataSize = cvImg.cols * cvImg.rows * sizeof(unsigned char);
checkCudaErrors(cudaMemcpy(dImgIn, cvImg.data, dataSize, cudaMemcpyHostToDevice));

// カーネル関数の実行
dim3 grid(cvImg.rows);
dim3 block(1024);
invertKernel << <grid, block >> >(dImgOut, dImgIn, cvImg.cols);

checkCudaErrors(cudaGraphicsUnmapResources(1, &cuda_pbo_resource, 0));

glBindTexture(GL_TEXTURE_2D, tex);
glBindBuffer(GL_PIXEL_UNPACK_BUFFER, pbo);
glTexSubImage2D(GL_TEXTURE_2D, 0, 0, 0, cvImg.cols, cvImg.rows,
GL_LUMINANCE, GL_UNSIGNED_BYTE, 0);
glBindBuffer(GL_PIXEL_UNPACK_BUFFER, 0);

glEnable(GL_TEXTURE_2D);
glBegin(GL_QUADS);
{
glTexCoord2f(0, 1);
glVertex2f(-1, -1);
glTexCoord2f(1, 1);
glVertex2f(1, -1);
glTexCoord2f(1, 0);
glVertex2f(1, 1);
glTexCoord2f(0, 0);
glVertex2f(-1, 1);
}
glEnd();

glBindTexture(GL_TEXTURE_2D, 0);
glutSwapBuffers();
}
display関数です。
PBOのポインタ取得、入力画像のGPU転送、カーネル実行、テクスチャマッピングを行っています。
処理結果をCPUに戻さずに、PBOからテクスチャにコピー(GPU内で閉じてる)して描画しているのがポイントです。

念のため、カーネル関数の一例も載せておきます。
__global__ void invertKernel(unsigned char *out, const unsigned char *in, const int width)
{
for (int i = threadIdx.x; i < width; i += blockDim.x) {
int index = blockIdx.x * width + i;
out[index] = 255 - in[index];
}
}

以上です。