Tested ok, pushed with the previous patch, Thanks.
On Mon, Nov 10, 2014 at 01:49:19PM +0800, Guo Yejun wrote: > Signed-off-by: Guo Yejun <[email protected]> > --- > kernels/runtime_climage_from_boname.cl | 8 ++ > utests/CMakeLists.txt | 11 +- > utests/runtime_climage_from_boname.cpp | 208 > +++++++++++++++++++++++++++++++++ > 3 files changed, 226 insertions(+), 1 deletion(-) > create mode 100644 kernels/runtime_climage_from_boname.cl > create mode 100644 utests/runtime_climage_from_boname.cpp > > diff --git a/kernels/runtime_climage_from_boname.cl > b/kernels/runtime_climage_from_boname.cl > new file mode 100644 > index 0000000..2e9da0f > --- /dev/null > +++ b/kernels/runtime_climage_from_boname.cl > @@ -0,0 +1,8 @@ > +__kernel void > +runtime_climage_from_boname(__write_only image2d_t dst) > +{ > + int2 coord; > + coord.x = (int)get_global_id(0); > + coord.y = (int)get_global_id(1); > + write_imagef(dst, coord, (float4)(0.34)); > +} > diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt > index f609cc2..405ee43 100644 > --- a/utests/CMakeLists.txt > +++ b/utests/CMakeLists.txt > @@ -201,6 +201,15 @@ set (utests_sources > utest_file_map.cpp > utest_helper.cpp) > > +if (X11_FOUND) > + set(utests_sources > + ${utests_sources} > + runtime_climage_from_boname.cpp) > + SET(UTESTS_REQUIRED_X11_LIB ${X11_LIBRARIES} ${XEXT_LIBRARIES}) > +else() > +SET(UTESTS_REQUIRED_X11_LIB "") > +endif (X11_FOUND) > + > SET (kernel_bin ${CMAKE_CURRENT_SOURCE_DIR}/../kernels/compiler_ceil) > > if(GEN_PCI_ID) > @@ -249,7 +258,7 @@ endif () > > ADD_LIBRARY(utests SHARED ${ADDMATHFUNC} ${utests_sources}) > > -TARGET_LINK_LIBRARIES(utests cl m ${OPENGL_LIBRARIES} > ${UTESTS_REQUIRED_EGL_LIB} ${CMAKE_THREAD_LIBS_INIT}) > +TARGET_LINK_LIBRARIES(utests cl m ${OPENGL_LIBRARIES} > ${UTESTS_REQUIRED_EGL_LIB} ${CMAKE_THREAD_LIBS_INIT} > ${UTESTS_REQUIRED_X11_LIB}) > > ADD_EXECUTABLE(utest_run utest_run.cpp) > TARGET_LINK_LIBRARIES(utest_run utests) > diff --git a/utests/runtime_climage_from_boname.cpp > b/utests/runtime_climage_from_boname.cpp > new file mode 100644 > index 0000000..30bbdbd > --- /dev/null > +++ b/utests/runtime_climage_from_boname.cpp > @@ -0,0 +1,208 @@ > +#include <cstdint> > +#include <cstring> > +#include <iostream> > +#include "utest_helper.hpp" > +#include "utest_file_map.hpp" > + > +#include <stdlib.h> > +#include <fcntl.h> > +#include <unistd.h> > + > +extern "C" > +{ > +#include <X11/Xlibint.h> > +#include <X11/Xlib.h> > +#include <xf86drm.h> > +#include <intel_bufmgr.h> > +#include <drm.h> > +#include <drm_sarea.h> > +#include <X11/Xmd.h> > +#include <X11/Xregion.h> > +#include <X11/extensions/Xext.h> > +#include <X11/extensions/extutil.h> > +} > + > +// part of following code is copy from beignet/src/x11/ > +typedef struct { > + CARD8 reqType; > + CARD8 dri2Reqtype; > + CARD16 length B16; > + CARD32 window B32; > + CARD32 magic B32; > +} xDRI2AuthenticateReq; > +#define sz_xDRI2AuthenticateReq 12 > + > +typedef struct { > + BYTE type; /* X_Reply */ > + BYTE pad1; > + CARD16 sequenceNumber B16; > + CARD32 length B32; > + CARD32 authenticated B32; > + CARD32 pad2 B32; > + CARD32 pad3 B32; > + CARD32 pad4 B32; > + CARD32 pad5 B32; > + CARD32 pad6 B32; > +} xDRI2AuthenticateReply; > +#define sz_xDRI2AuthenticateReply 32 > + > +#define X_DRI2Authenticate 2 > + > +static char va_dri2ExtensionName[] = "DRI2"; > +static XExtensionInfo _va_dri2_info_data; > +static XExtensionInfo *va_dri2Info = &_va_dri2_info_data; > +static XEXT_GENERATE_CLOSE_DISPLAY (VA_DRI2CloseDisplay, va_dri2Info) > +static /* const */ XExtensionHooks va_dri2ExtensionHooks = { > + NULL, /* create_gc */ > + NULL, /* copy_gc */ > + NULL, /* flush_gc */ > + NULL, /* free_gc */ > + NULL, /* create_font */ > + NULL, /* free_font */ > + VA_DRI2CloseDisplay, /* close_display */ > + NULL, /* wire_to_event */ > + NULL, /* event_to_wire */ > + NULL, /* error */ > + NULL, /* error_string */ > +}; > + > +static XEXT_GENERATE_FIND_DISPLAY (DRI2FindDisplay, va_dri2Info, > + va_dri2ExtensionName, > + &va_dri2ExtensionHooks, > + 0, NULL) > + > +static Bool VA_DRI2Authenticate(Display *dpy, XID window, drm_magic_t magic) > +{ > + XExtDisplayInfo *info = DRI2FindDisplay(dpy); > + xDRI2AuthenticateReq *req; > + xDRI2AuthenticateReply rep; > + > + XextCheckExtension (dpy, info, va_dri2ExtensionName, False); > + > + LockDisplay(dpy); > + GetReq(DRI2Authenticate, req); > + req->reqType = info->codes->major_opcode; > + req->dri2Reqtype = X_DRI2Authenticate; > + req->window = window; > + req->magic = magic; > + > + if (!_XReply(dpy, (xReply *)&rep, 0, xFalse)) { > + UnlockDisplay(dpy); > + SyncHandle(); > + return False; > + } > + > + UnlockDisplay(dpy); > + SyncHandle(); > + > + return rep.authenticated; > +} > + > + > +void runtime_climage_from_boname(void) > +{ > + const int w = 1024; > + const int h = 256; > + const int hStart = 128; > + const int offset = hStart * w; > + > + // Setup kernel and buffers > + OCL_CREATE_KERNEL("runtime_climage_from_boname"); > + > + int fd = open("/dev/dri/card0", O_RDWR); > + OCL_ASSERT(fd>0); > + > + drm_magic_t magic; > + drmGetMagic(fd, &magic); > + > + Display* dpy = XOpenDisplay(NULL); > + XID root = RootWindow(dpy, DefaultScreen(dpy)); > + > + Bool auth = VA_DRI2Authenticate(dpy, root, magic); > + OCL_ASSERT(auth); > + > + drm_intel_bufmgr* bufmgr = drm_intel_bufmgr_gem_init(fd, 1024); > + OCL_ASSERT(bufmgr != NULL); > + > + drm_intel_bo * bo = drm_intel_bo_alloc(bufmgr, > "runtime_climage_from_boname", w*h, 0); > + OCL_ASSERT(bo != NULL); > + > + drm_intel_bo_map(bo, 0); > + unsigned char* addr = (unsigned char*)bo->virt; > + memset(addr, 0xCD, w*h); > + drm_intel_bo_unmap(bo); > + > + unsigned int boName = 0; > + drm_intel_bo_flink(bo, &boName); > + > + cl_image_format fmt; > + fmt.image_channel_order = CL_R; > + fmt.image_channel_data_type = CL_UNORM_INT8; > + > + cl_libva_image imageParam; > + imageParam.fmt = fmt; > + imageParam.bo_name = boName; > + imageParam.offset = offset; > + imageParam.width = w; > + imageParam.height = h - hStart; > + imageParam.row_pitch = w; > + > + cl_mem dst = clCreateImageFromLibvaIntel(ctx, &imageParam, NULL); > + > + // Run the kernel > + OCL_SET_ARG(0, sizeof(cl_mem), &dst); > + globals[0] = w; > + globals[1] = h-hStart; > + locals[0] = 16; > + locals[1] = 16; > + OCL_NDRANGE(2); > + > + OCL_FINISH(); > + > + drm_intel_bo_map(bo, 0); > + addr = (unsigned char*)bo->virt; > + for (int i = 0; i < hStart; ++i) { > + for (int j = 0; j < w; ++j) { > + OCL_ASSERT(addr[j+i*w]==0xCD); > + } > + } > + for (int i = hStart; i < h; ++i) { > + for (int j = 0; j < w; ++j) { > + OCL_ASSERT(addr[j+i*w]==(unsigned char)(0.34*255+0.5)); > + } > + } > + drm_intel_bo_unmap(bo); > + > + > + // Run the kernel for the seconde time > + OCL_SET_ARG(0, sizeof(cl_mem), &dst); > + globals[0] = w; > + globals[1] = h-hStart; > + locals[0] = 16; > + locals[1] = 16; > + OCL_NDRANGE(2); > + > + OCL_FINISH(); > + > + drm_intel_bo_map(bo, 0); > + addr = (unsigned char*)bo->virt; > + for (int i = 0; i < hStart; ++i) { > + for (int j = 0; j < w; ++j) { > + OCL_ASSERT(addr[j+i*w]==0xCD); > + } > + } > + for (int i = hStart; i < h; ++i) { > + for (int j = 0; j < w; ++j) { > + OCL_ASSERT(addr[j+i*w]==(unsigned char)(0.34*255+0.5)); > + } > + } > + drm_intel_bo_unmap(bo); > + > + clReleaseMemObject(dst); > + drm_intel_bo_unreference(bo); > + drm_intel_bufmgr_destroy(bufmgr); > + XCloseDisplay(dpy); > + close(fd); > +} > + > +MAKE_UTEST_FROM_FUNCTION(runtime_climage_from_boname); > -- > 2.1.0 > > _______________________________________________ > Beignet mailing list > [email protected] > http://lists.freedesktop.org/mailman/listinfo/beignet _______________________________________________ Beignet mailing list [email protected] http://lists.freedesktop.org/mailman/listinfo/beignet
