static void _update_pf_viewer(QSP_ARG_DECL Platform_Viewer *pvp, Data_Obj *dp)
{
#ifdef HAVE_OPENGL
int t;
//cudaError_t e;
// unmap buffer before using w/ GL
if( BUF_IS_MAPPED(dp) ){
if( (*PF_UNMAPBUF_FN(PFDEV_PLATFORM(OBJ_PFDEV(dp))))
(QSP_ARG dp) < 0 ) {
warn("update_pf_viewer: buffer unmap error!?");
}
CLEAR_OBJ_FLAG_BITS(dp, DT_BUF_MAPPED);
// propagate change to children and parents
propagate_flag(dp,DT_BUF_MAPPED);
}
glClear(GL_COLOR_BUFFER_BIT);
glBindTexture(GL_TEXTURE_2D, OBJ_TEX_ID(dp));
// is glBindBuffer REALLY part of libGLEW???
//#ifdef HAVE_LIBGLEW
glBindBuffer(GL_PIXEL_UNPACK_BUFFER, OBJ_BUF_ID(dp));
//#endif // HAVE_LIBGLEW
t=gl_pixel_type(dp);
glTexSubImage2D(GL_TEXTURE_2D, 0, // target, level
0, 0, // x0, y0
OBJ_COLS(dp), OBJ_ROWS(dp), // dx, dy
t,
GL_UNSIGNED_BYTE, // type
OFFSET(0)); // offset into PIXEL_UNPACK_BUFFER
//#ifdef HAVE_LIBGLEW
glBindBuffer(GL_PIXEL_UNPACK_BUFFER, 0);
//#endif // HAVE_LIBGLEW
glBegin(GL_QUADS);
glTexCoord2f(0, 1); glVertex2f(-1.0, -1.0);
glTexCoord2f(0, 0); glVertex2f(-1.0, 1.0);
glTexCoord2f(1, 0); glVertex2f(1.0, 1.0);
glTexCoord2f(1, 1); glVertex2f(1.0, -1.0);
glEnd();
glBindTexture(GL_TEXTURE_2D, 0);
if( (*PF_MAPBUF_FN(PFDEV_PLATFORM(OBJ_PFDEV(dp))))(QSP_ARG dp) < 0 ){
warn("update_pf_viewer: Error mapping buffer!?");
}
SET_OBJ_FLAG_BITS(dp, DT_BUF_MAPPED);
// propagate change to children and parents
propagate_flag(dp,DT_BUF_MAPPED);
#else // ! HAVE_OPENGL
NO_OGL_MSG
#endif // ! HAVE_OPENGL
}
static void ocl_run_kernel(QSP_ARG_DECL void *kp, Vec_Expr_Node *arg_enp, Platform_Device *pdp)
{
cl_kernel kernel;
cl_int status;
cl_event event;
int karg_idx=0;
size_t global_work_size[3] = {1, 1, 1};
kernel = kp;
global_work_size[0] = set_fused_kernel_args(kernel, &karg_idx, arg_enp, PFDEV_PLATFORM(pdp));
//fprintf(stderr,"ocl_run_kernel: global work size = %ld\n",global_work_size[0]);
status = clEnqueueNDRangeKernel(
OCLDEV_QUEUE( pdp ),
kernel,
1, /* work_dim, 1-3 */
NULL,
global_work_size,
/*local_work_size*/ NULL,
0, /* num_events_in_wait_list */
NULL, /* event_wait_list */
&event /* event */
);
if( status != CL_SUCCESS )
report_ocl_error(status, "clEnqueueNDRangeKernel" );
clWaitForEvents(1,&event);
}
void show_gpu_vector(QSP_ARG_DECL Platform_Device *pdp, void *ptr, int len )
{
// BUG we assume float type!?
float *buf;
size_t siz;
int i;
siz= len*sizeof(float);
buf=malloc(siz);
if( buf==NULL ) NERROR1("show_gpu_vector: error allocating buffer!?");
fprintf(stderr,"show_gpu_vector: src = 0x%lx\n",(long)ptr);
// now do the memory transfer
(*PF_MEM_DNLOAD_FN(PFDEV_PLATFORM(pdp)))(QSP_ARG buf, ptr, siz, 0, pdp );
for(i=0;i<len;i++){
fprintf(stderr,"%d\t%g\n",i,buf[i]);
}
free(buf);
}
static void init_ocl_dev_memory(QSP_ARG_DECL Platform_Device *pdp)
{
char area_name[MAX_AREA_NAME_LEN+1];
Data_Area *ap;
//strcpy(area_name,PFDEV_NAME(pdp));
// make sure names will fit - longest name is %s.%s_host_mapped
if( strlen(PLATFORM_NAME(PFDEV_PLATFORM(pdp)))+strlen(PFDEV_NAME(pdp))+strlen("._host_mapped") > MAX_AREA_NAME_LEN )
error1("init_ocl_dev_memory: area name too large for buffer, increase MAX_AREA_NAME_LEN!?");
sprintf(area_name,"%s.%s",
PLATFORM_NAME(PFDEV_PLATFORM(pdp)),PFDEV_NAME(pdp));
// what should the name for the memory area be???
// address set to NULL says use custom allocator - see dobj/makedobj.c
ap = pf_area_init(area_name,NULL,0, MAX_OCL_GLOBAL_OBJECTS,DA_OCL_GLOBAL,pdp);
if( ap == NULL ){
sprintf(ERROR_STRING,
"init_ocl_dev_memory: error creating global data area %s",area_name);
warn(ERROR_STRING);
}
// g++ won't take this line!?
SET_AREA_PFDEV(ap,pdp);
// BUG should be per-device, not global table...
pdp->pd_ap[PF_GLOBAL_AREA_INDEX] = ap;
/* We used to declare a heap for constant memory here,
* but there wasn't much of a point because:
* Constant memory can't be allocated, rather it is declared
* in the .cu code, and placed by the compiler as it sees fit.
* To have objects use this, we would have to declare a heap and
* manage it ourselves...
* There's only 64k, so we should be sparing...
* We'll try this later...
*/
/* Make up another area for the host memory
* which is locked and mappable to the device.
* We don't allocate a pool here, but do it as needed...
*/
//strcat(cname,"_host");
sprintf(area_name,"%s.%s_host",
PLATFORM_NAME(PFDEV_PLATFORM(pdp)),PFDEV_NAME(pdp));
ap = pf_area_init(area_name,(u_char *)NULL,0,MAX_OCL_MAPPED_OBJECTS,
DA_OCL_HOST,pdp);
if( ap == NULL ){
sprintf(ERROR_STRING,
"init_ocl_dev_memory: error creating host data area %s",area_name);
error1(ERROR_STRING);
}
SET_AREA_PFDEV(ap, pdp);
pdp->pd_ap[PF_HOST_AREA_INDEX] = ap;
/* Make up another psuedo-area for the mapped host memory;
* This is the same memory as above, but mapped to the device.
* In the current implementation, we create objects in the host
* area, and then automatically create an alias on the device side.
* There is a BUG in that by having this psuedo area in the data
* area name space, a user could select it as the data area and
* then try to create an object. We will detect this in make_dobj,
* and complain.
*/
//strcpy(cname,dname);
//strcat(cname,"_host_mapped");
sprintf(area_name,"%s.%s_host_mapped",
PLATFORM_NAME(PFDEV_PLATFORM(pdp)),PFDEV_NAME(pdp));
ap = pf_area_init(area_name,(u_char *)NULL,0,MAX_OCL_MAPPED_OBJECTS,
DA_OCL_HOST_MAPPED,pdp);
if( ap == NULL ){
sprintf(ERROR_STRING,
"init_ocl_dev_memory: error creating host-mapped data area %s",area_name);
error1(ERROR_STRING);
}
SET_AREA_PFDEV(ap,pdp);
pdp->pd_ap[PF_HOST_MAPPED_AREA_INDEX] = ap;
if( verbose ){
sprintf(ERROR_STRING,"init_ocl_dev_memory DONE");
advise(ERROR_STRING);
}
}
// This is the normal display path
static void update_pf_viewer(QSP_ARG_DECL Platform_Viewer *pvp, Data_Obj *dp)
{
#ifdef HAVE_OPENGL
int t;
//cudaError_t e;
// unmap buffer before using w/ GL
if( BUF_IS_MAPPED(dp) ){
if( (*PF_UNMAPBUF_FN(PFDEV_PLATFORM(OBJ_PFDEV(dp))))
(QSP_ARG dp) < 0 ) {
WARN("update_pf_viewer: buffer unmap error!?");
}
#ifdef FOOBAR
e = cudaGLUnmapBufferObject( OBJ_BUF_ID(dp) );
if( e != cudaSuccess ){
describe_cuda_driver_error2("update_pf_viewer",
"cudaGLUnmapBufferObject",e);
NERROR1("failed to unmap buffer object");
}
#endif // FOOBAR
CLEAR_OBJ_FLAG_BITS(dp, DT_BUF_MAPPED);
// propagate change to children and parents
propagate_flag(dp,DT_BUF_MAPPED);
}
//
//bind_texture(OBJ_DATA_PTR(dp));
glClear(GL_COLOR_BUFFER_BIT);
/*
sprintf(ERROR_STRING,"update_pf_viewer: tex_id = %d, buf_id = %d",
OBJ_TEX_ID(dp),OBJ_BUF_ID(dp));
advise(ERROR_STRING);
*/
glBindTexture(GL_TEXTURE_2D, OBJ_TEX_ID(dp));
// is glBindBuffer REALLY part of libGLEW???
//#ifdef HAVE_LIBGLEW
glBindBuffer(GL_PIXEL_UNPACK_BUFFER, OBJ_BUF_ID(dp));
//#endif // HAVE_LIBGLEW
#ifdef FOOBAR
switch(OBJ_COMPS(dp)){
/* what used to be here??? */
}
#endif /* FOOBAR */
t=gl_pixel_type(dp);
glTexSubImage2D(GL_TEXTURE_2D, 0, // target, level
0, 0, // x0, y0
OBJ_COLS(dp), OBJ_ROWS(dp), // dx, dy
t,
GL_UNSIGNED_BYTE, // type
OFFSET(0)); // offset into PIXEL_UNPACK_BUFFER
//#ifdef HAVE_LIBGLEW
glBindBuffer(GL_PIXEL_UNPACK_BUFFER, 0);
//#endif // HAVE_LIBGLEW
glBegin(GL_QUADS);
glTexCoord2f(0, 1); glVertex2f(-1.0, -1.0);
glTexCoord2f(0, 0); glVertex2f(-1.0, 1.0);
glTexCoord2f(1, 0); glVertex2f(1.0, 1.0);
glTexCoord2f(1, 1); glVertex2f(1.0, -1.0);
glEnd();
glBindTexture(GL_TEXTURE_2D, 0);
#ifdef FOOBAR
e = cudaGLMapBufferObject( &OBJ_DATA_PTR(dp), OBJ_BUF_ID(dp) );
if( e != cudaSuccess ){
WARN("Error mapping buffer object!?");
// should we return now, with possibly other cleanup???
}
#endif // FOOBAR
if( (*PF_MAPBUF_FN(PFDEV_PLATFORM(OBJ_PFDEV(dp))))(QSP_ARG dp) < 0 ){
WARN("update_pf_viewer: Error mapping buffer!?");
}
SET_OBJ_FLAG_BITS(dp, DT_BUF_MAPPED);
// propagate change to children and parents
propagate_flag(dp,DT_BUF_MAPPED);
#else // ! HAVE_OPENGL
NO_OGL_MSG
#endif // ! HAVE_OPENGL
}
static COMMAND_FUNC( do_new_gl_buffer )
{
const char *s;
Data_Obj *dp;
Platform_Device *pdp;
Compute_Platform *cdp;
dimension_t d,w,h;
#ifdef HAVE_OPENGL
Dimension_Set ds;
int t;
#endif // HAVE_OPENGL
s = NAMEOF("name for GL buffer object");
cdp = pick_platform("platform");
if( cdp != NULL )
push_pfdev_context(QSP_ARG PF_CONTEXT(cdp) );
pdp = pick_pfdev("device");
if( cdp != NULL )
pop_pfdev_context(SINGLE_QSP_ARG);
w = (int)HOW_MANY("width");
h = (int)HOW_MANY("height");
d = (int)HOW_MANY("depth");
/* what should the depth be??? default to 1 for now... */
if( pdp == NULL ) return;
/* Make sure this name isn't already in use... */
dp = dobj_of(s);
if( dp != NULL ){
sprintf(ERROR_STRING,"Data object name '%s' is already in use, can't use for GL buffer object.",s);
warn(ERROR_STRING);
return;
}
#ifdef HAVE_OPENGL
// BUG need to be able to set the cuda device.
// Note, however, that we don't need GL buffers on the Tesla...
//set_data_area(cuda_data_area[0][0]);
set_data_area( PFDEV_AREA(pdp,PFDEV_GLOBAL_AREA_INDEX) );
ds.ds_dimension[0]=d;
ds.ds_dimension[1]=w;
ds.ds_dimension[2]=h;
ds.ds_dimension[3]=1;
ds.ds_dimension[4]=1;
dp = _make_dp(QSP_ARG s,&ds,PREC_FOR_CODE(PREC_UBY));
if( dp == NULL ){
sprintf(ERROR_STRING,
"Error creating data_obj header for %s",s);
error1(ERROR_STRING);
}
SET_OBJ_FLAG_BITS(dp, DT_NO_DATA); /* can't free this data */
SET_OBJ_FLAG_BITS(dp, DT_GL_BUF); /* indicate obj is a GL buffer */
SET_OBJ_DATA_PTR(dp, NULL);
//fprintf(stderr,"do_new_gl_buffer: allocating gl_info for %s\n",OBJ_NAME(dp));
SET_OBJ_GL_INFO(dp, (GL_Info *) getbuf( sizeof(GL_Info) ) );
//fprintf(stderr,"do_new_gl_buffer: DONE allocating gl_info for %s\n",OBJ_NAME(dp));
glew_check(SINGLE_QSP_ARG); /* without this, we get a segmentation
* violation on glGenBuffers???
*/
// We need an extra field in which to store the GL identifier...
// AND another extra field in which to store the associated texid.
// Why is this ifdef here? These don't seem to depend
// on libglew???
// Answer: We need libglew to bring in openGL extensions like glBindBuffer...
//advise("calling glGenBuffers");
//fprintf(stderr,"OBJ_GL_INFO(%s) = 0x%lx\n",OBJ_NAME(dp),(long)OBJ_GL_INFO(dp));
//fprintf(stderr,"OBJ_BUF_ID_P(%s) = 0x%lx\n",OBJ_NAME(dp),(long)OBJ_BUF_ID_P(dp));
// BUG glGenBuffers seems to require v1.5???
glGenBuffers(1, OBJ_BUF_ID_P(dp) ); // first arg is # buffers to generate?
//sprintf(ERROR_STRING,"glGenBuffers gave us buf_id = %d",OBJ_BUF_ID(dp));
//advise(ERROR_STRING);
glBindBuffer(GL_PIXEL_UNPACK_BUFFER, OBJ_BUF_ID(dp) );
// glBufferData will allocate the memory for the buffer,
// but won't copy unless the pointer is non-null
// How do we get the gpu memory space address?
// That must be with map
glBufferData(GL_PIXEL_UNPACK_BUFFER,
OBJ_COMPS(dp) * OBJ_COLS(dp) * OBJ_ROWS(dp), NULL, GL_STREAM_DRAW);
/* buffer arg set to 0 unbinds any previously bound buffers...
* and restores client memory usage.
*/
glBindBuffer(GL_PIXEL_UNPACK_BUFFER, 0);
//#endif // HAVE_LIBGLEW
glGenTextures(1, OBJ_TEX_ID_P(dp) ); // makes a texture name
fprintf(stderr,"new_gl_buffer: new texture name is 0x%x\n",OBJ_TEX_ID(dp));
glBindTexture(GL_TEXTURE_2D, OBJ_TEX_ID(dp) );
t = gl_pixel_type(dp);
glTexImage2D( GL_TEXTURE_2D,
0, // level-of-detail - is this the same as miplevel???
OBJ_COMPS(dp), // internal format, can also be symbolic constant such as
// GL_RGBA etc
OBJ_COLS(dp), // width - must be 2^n+2 (border) for some n???
OBJ_ROWS(dp), // height - must be 2^m+2 (border) for some m???
0, // border - must be 0 or 1
t, // format of pixel data
GL_UNSIGNED_BYTE, // type of pixel data
NULL // pixel data - null pointer means
// allocate but do not copy?
// - offset into PIXEL_UNPACK_BUFFER??
);
glTexParameterf(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_LINEAR);
glTexParameterf(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_LINEAR);
// Why was this here? It would seem to un-bind the target???
glBindTexture(GL_TEXTURE_2D, 0);
//glFinish(); // necessary or not?
//advise("calling platform-specific buffer registration function");
if( (*PF_REGBUF_FN(PFDEV_PLATFORM(pdp)))( QSP_ARG dp ) < 0 ){
WARN("do_new_gl_buffer: Error in platform-specific buffer registration!?");
// BUG? - should clean up here!
}
// Leave the buffer mapped by default
//cutilSafeCall(cudaGLMapBufferObject( &OBJ_DATA_PTR(dp), OBJ_BUF_ID(dp) ));
//advise("calling platform-specific buffer mapping function");
if( (*PF_MAPBUF_FN(PFDEV_PLATFORM(pdp)))( QSP_ARG dp ) < 0 ){
WARN("do_new_gl_buffer: Error in platform-specific buffer mapping!?");
// BUG? - should clean up here!
}
SET_OBJ_FLAG_BITS(dp, DT_BUF_MAPPED);
// propagate change to children and parents
propagate_flag(dp,DT_BUF_MAPPED);
#else // ! HAVE_OPENGL
NO_OGL_MSG
#endif // ! HAVE_OPENGL
} /* end do_new_gl_buffer */
