OpenGLでテクスチャを設定するとき，一番簡単な方法としては図のようなアプローチをとります．
この方法は，GPU側で用意したテクスチャメモリに向けて，CPUのメモリのデータをコピーします．
テクスチャをいくつか用意して，それを次々と切り替えたいときは，PBO（Pixel Buffer Object）図のようなアプローチもあります．
PBOからテクスチャへのメモリコピーは，GPU内だけで高速に行えるので，テクスチャの書き換えの都合に合わせてCPUの挙動を考慮する必要がなくなります．
ただ，この方法はいずれにしてもどこかの時点でCPUからGPUへのメモリコピーを行わなければなりません．
現在のコンピュータは，PCIeの通信もそこそこ高速になってきていますので，解像度が高すぎなければ，頻繁にCPUからGPUへのメモリコピーを行なっても問題は少ないかもしれません．（私は，毎フレームglTexImage2DでCPUからGPUのテクスチャメモリへコピーして，動画をテクスチャとして表示するというような実装をよくやりますが，解像度がそんなに高くなければ全く問題ありません）
しかし，CUDAやOpenCLで何かしらの計算をして，その計算結果を濃淡画像として表示する処理をOpenGLのテクスチャを使って実現すると，図のような方法になってしまいます．
データがGPU-CPU間を1往復しています．もしこれが，CPUメモリのデータを入力として使う処理の場合（例えば，画像を色変換して表示する），更にもう一回CPUからGPUへのコピーが増えて，1.5往復することになってしまいます．多くの場合，GPUで計算する際の計算そのものより，CPUからGPUへのデータコピーの方がよほど時間がかかります．
これではせっかくのGPU計算なのに非常に効率が悪いので，こうした，GPU計算の可視化のために，例えばCUDAでは，PBOとして確保した領域をCUDAカーネル内で操作することができます．
この方法をとることで，CPUからGPUへのコピーをすることなく，テクスチャの表示を行うことができます．CPUメモリにあるデータをGPUで処理して，可視化を行う場合でも，CPU-GPU間のデータコピーは1回で済みます．
CUDAを使ったその方法は，この本に詳しくでています．
GPU 並列図形処理入門 ――CUDA・OpenGLの導入と活用
こうしたGPU計算とPBOの連動は，OpenCLでもできます．しかし，多くの資料では「できまっせ〜」みたいな感じのことが書いてあるだけで，具体的にどうやるのかをちゃんと示していません．なので，色々な情報をかき集めて，やってみました．
GPUでFDTD法を計算し，その結果を濃淡画像としてテクスチャにします．
重要なのは以下の部分．
//OpenCLの計算に使え，尚且つOpenGLテクスチャとして表示可能なメモリを確保（サイズはglBufferDataARBで設定したサイズで確保される）
d_g_data=clCreateFromGLBuffer(context, CL_MEM_READ_WRITE, pbo, &ret);
そして以下．
clEnqueueAcquireGLObjects(command_queue,1,&d_g_data,0,NULL,NULL);//テクスチャ描画可能なメモリであることを通知
全体としてはこのようになります．このソースの中で，CPUからGPUへのメモリコピーが一回も発生していないことに注目してください．
#include<stdio.h> #include<stdlib.h> #include<string.h> #include"glew.h" #include<GLUT/GLUT.h> #include<OpenCL/cl.h> #include<OpenCL/cl_gl.h> #include <OpenCL/cl_gl_ext.h>// for CL_CONTEXT_PROPERTY_USE_CGL_SHAREGROUP_APPLE #include <OpenGL/CGLDevice.h> //for CGLGetShareGroup() #include <OpenGL/CGLCurrent.h> #define WIDTH 512 #define HEIGHT 512 #define MAX_PLATFORMS (10) #define MAX_DEVICES (10) void disp(); GLuint pbo; GLuint tex; cl_mem sf=NULL; //現在の値 cl_mem sf_1=NULL; //1回前の値 cl_mem sf_next=NULL; //次の値 cl_mem d_g_data=NULL; //テクスチャとして表示されるメモリ cl_command_queue command_queue = NULL; cl_kernel initval; //初期値設定カーネル cl_kernel fdtd; //FDTD実行カーネル cl_kernel texcopy; //計算結果をd_g_dataにコピーするカーネル size_t local_item_size[2]={16,16}; size_t global_item_size[2]={WIDTH,HEIGHT}; int width=WIDTH,height=HEIGHT; int main(int argc,char *argv[]){ glutInit(&argc,argv); glutInitDisplayMode(GLUT_RGBA|GLUT_DOUBLE); glutInitWindowSize(WIDTH,HEIGHT); glutCreateWindow("render"); glutDisplayFunc(disp); glewInit(); glEnable(GL_TEXTURE_2D); //PBOを設定 glGenBuffersARB(1,&pbo); glBindBufferARB(GL_PIXEL_UNPACK_BUFFER_ARB,pbo); glBufferDataARB(GL_PIXEL_UNPACK_BUFFER_ARB,WIDTH*HEIGHT*3*sizeof(unsigned char),0,GL_STREAM_DRAW_ARB); glBindBufferARB(GL_PIXEL_UNPACK_BUFFER_ARB,0); //テクスチャを設定 glGenTextures(1,&tex); glBindTexture(GL_TEXTURE_2D,tex); glTexImage2D(GL_TEXTURE_2D,0,GL_RGB,WIDTH,HEIGHT,0,GL_RGB,GL_UNSIGNED_BYTE,NULL); glTexParameteri(GL_TEXTURE_2D,GL_TEXTURE_MAG_FILTER,GL_LINEAR); glTexParameteri(GL_TEXTURE_2D,GL_TEXTURE_MIN_FILTER,GL_LINEAR); glBindTexture(GL_TEXTURE_2D,0); //OpenCL設定 FILE *fp; cl_platform_id platformid[MAX_PLATFORMS]; cl_uint numplatforms; cl_uint status = 1; cl_uint platform = 0; int device = 0; cl_device_id device_id[MAX_DEVICES]; cl_uint num_devices; char str[BUFSIZ]; size_t ret_size; cl_int ret; cl_context context = NULL; cl_program program = NULL; char *source_str= new char[100000];//(char *)malloc(100000 * sizeof(char)); size_t source_size; clGetPlatformIDs(MAX_PLATFORMS, platformid, &numplatforms); clGetDeviceIDs(platformid[platform], CL_DEVICE_TYPE_ALL, MAX_DEVICES, device_id, &num_devices); for(device=0;device<num_devices;device++){ clGetDeviceInfo(device_id[device], CL_DEVICE_NAME, sizeof(str), str, &ret_size); printf("%s %d %d\n",str,device,device_id[device]); if(strcmp(str,"AMD Radeon Pro 560 Compute Engine")==0) break; //"AMD Radeon Pro 560 Compute Engine"の番号のdeviceをオープンする（そのために"AMD Radeon Pro 560 Compute Engine"を見つけた時点でループを抜ける） } //Mac OSではこの部分が必要（Windowsでは不要だった） CGLContextObj glContext = CGLGetCurrentContext(); CGLShareGroupObj shareGroup = CGLGetShareGroup(glContext); cl_context_properties props[] = { CL_CONTEXT_PROPERTY_USE_CGL_SHAREGROUP_APPLE, (cl_context_properties)shareGroup, 0 }; //Mac OSではこの部分までが必要（Windowsでは不要だった） context = clCreateContext(props, 1, &device_id[device], NULL, NULL, &ret); command_queue = clCreateCommandQueue(context, device_id[device], 0, &ret); if ((fp = fopen("kernel.cl", "r")) == NULL) { fprintf(stderr, "kernel source open error\n"); getchar(); exit(1); } source_size = fread(source_str, sizeof(char), 100000, fp); fclose(fp); program = clCreateProgramWithSource(context, 1, (const char **)&source_str, (const size_t *)&source_size, &ret); //OpenCLプログラムのビルド clBuildProgram(program, 1, &device_id[device], NULL, NULL, NULL); //カーネルの設定 texcopy=clCreateKernel(program, "texcopy", &ret); fdtd=clCreateKernel(program, "fdtd", &ret); initval=clCreateKernel(program, "initval", &ret); //GPU上のメモリを確保 sf=clCreateBuffer(context, CL_MEM_READ_WRITE, WIDTH*HEIGHT*sizeof(float), NULL, &ret); sf_1=clCreateBuffer(context, CL_MEM_READ_WRITE, WIDTH*HEIGHT*sizeof(float), NULL, &ret); sf_next=clCreateBuffer(context, CL_MEM_READ_WRITE, WIDTH*HEIGHT*sizeof(float), NULL, &ret); //OpenCLの計算に使え，尚且つOpenGLテクスチャとして表示可能なメモリを確保（サイズはglBufferDataARBで設定したサイズで確保される） d_g_data=clCreateFromGLBuffer(context, CL_MEM_READ_WRITE, pbo, &ret); //初期化カーネル実行 clSetKernelArg(initval, 0, sizeof(cl_mem), &sf); clSetKernelArg(initval, 1, sizeof(cl_mem), &sf_1); clSetKernelArg(initval, 2, sizeof(int), &width); clSetKernelArg(initval, 3, sizeof(int), &height); clEnqueueNDRangeKernel(command_queue, initval, 2, NULL, global_item_size, local_item_size, 0, NULL, NULL); //レンダリング開始 glutMainLoop(); clReleaseMemObject(sf); clReleaseMemObject(sf_1); clReleaseMemObject(sf_next); return 0; } //レンダリング関数 void disp(){ //FDTD実行 clSetKernelArg(fdtd, 0, sizeof(cl_mem), &sf); clSetKernelArg(fdtd, 1, sizeof(cl_mem), &sf_1); clSetKernelArg(fdtd, 2, sizeof(cl_mem), &sf_next); clSetKernelArg(fdtd, 3, sizeof(int), &width); clSetKernelArg(fdtd, 4, sizeof(int), &height); clEnqueueNDRangeKernel(command_queue, fdtd, 2, NULL, global_item_size, local_item_size, 0, NULL, NULL); //テクスチャ用メモリへコピー clEnqueueAcquireGLObjects(command_queue,1,&d_g_data,0,NULL,NULL);//テクスチャ描画可能なメモリであることを通知 clSetKernelArg(texcopy,0,sizeof(cl_mem),&d_g_data); clSetKernelArg(texcopy,1,sizeof(cl_mem),&sf); clSetKernelArg(texcopy,2,sizeof(cl_mem),&sf_1); clSetKernelArg(texcopy,3,sizeof(cl_mem),&sf_next); clSetKernelArg(texcopy, 4, sizeof(int), &width); clSetKernelArg(texcopy, 5, sizeof(int), &height); clEnqueueNDRangeKernel(command_queue,texcopy,2,NULL,global_item_size,local_item_size,0,NULL,NULL); clEnqueueReleaseGLObjects(command_queue,1,&d_g_data,0,NULL,NULL);//毎回解放する必要がある //描画（CPUからGPUへのメモリコピーが発生していないことに注目） glBindBuffer(GL_PIXEL_UNPACK_BUFFER,pbo); glBindTexture(GL_TEXTURE_2D,tex); glTexSubImage2D(GL_TEXTURE_2D,0,0,0,WIDTH,HEIGHT,GL_RGB,GL_UNSIGNED_BYTE,NULL); glBindTexture(GL_TEXTURE_2D,0); glBindBuffer(GL_PIXEL_UNPACK_BUFFER,0); glClear(GL_COLOR_BUFFER_BIT|GL_DEPTH_BUFFER_BIT); glBindTexture(GL_TEXTURE_2D,tex); glTexEnvi(GL_TEXTURE_ENV, GL_TEXTURE_ENV_MODE, GL_REPLACE); glBegin(GL_QUADS); glTexCoord2f(0,0); glVertex3d(-1, -1, 0); glTexCoord2f(0,1); glVertex3d(-1, 1, 0); glTexCoord2f(1,1); glVertex3d(1, 1, 0); glTexCoord2f(1,0); glVertex3d(1, -1, 0); glEnd(); glFlush(); glBindTexture(GL_TEXTURE_2D,0); glutSwapBuffers(); glutPostRedisplay(); }
//初期値設定
__kernel void initval(__global float *sf,__global float *sf_1,int width,int height){
int idx=get_global_id(0);
int idy=get_global_id(1);
int index;
float pi=3.14159265358979323;
index=idx+(idy*width);
sf[index]=0;
sf_1[index]=0;
if((idx>=width/2)&&(idx<width/2+40)&&(idy>=width/2)&&(idy<width/2+40)){
sf_1[index]=0.5-0.5*cos(2*pi*(float)(idx-width/2)/40);
sf_1[index]*=0.5-0.5*cos(2*pi*(float)(idy-width/2)/40);
sf_1[index]*=0.73;
}
}
//FDTD法
__kernel void fdtd(__global float *sf,__global float *sf_1,__global float *sf_next,int width,int height){
int idx=get_global_id(0);
int idy=get_global_id(1);
int index=idx+(idy*width);
int index1=idx+((idy-1)*width);
int index2=idx+((idy+1)*width);
int index3=(idx-1)+(idy*width);
int index4=(idx+1)+(idy*width);
float deltat=0.00001,deltax=0.001,c=70.0f;
sf_next[index]=-(sf_1[index]-2*sf[index])+(c*c*deltat*deltat*(sf[index4]+sf[index3]-2*sf[index])/(deltax*deltax))+(c*c*deltat*deltat*(sf[index2]+sf[index1]-2*sf[index])/(deltax*deltax));
}
//計算結果をテクスチャ用メモリへコピー
__kernel void texcopy(__global unsigned char *tex,__global float *sf,__global float *sf_1,__global float *sf_next,int width,int height){
int idx=get_global_id(0);
int idy=get_global_id(1);
int index=idx+(idy*width);
int tindex=index*3;
sf_1[index]=sf[index];
sf[index]=sf_next[index];
//texはsfの3倍のメモリサイズなので（RGB画像だから）3ピクセルごと同じ値をコピーする
tex[tindex]=(unsigned char)(sf[index]*255.0f);
tex[tindex+1]=(unsigned char)(sf[index]*255.0f);
tex[tindex+2]=(unsigned char)(sf[index]*255.0f);
}
OpenCLの初期化のところで，動作させるOSによって処理内容が変わるようです．本記事のコードは，MacOSX 10.13.4，Xcode 9.3でのコードです．
