<div dir="ltr"><br><div class="gmail_extra"><br><br><div class="gmail_quote">On Tue, Sep 17, 2013 at 1:42 PM, Wenju He <span dir="ltr"><<a href="mailto:wenju@multicorewareinc.com" target="_blank">wenju@multicorewareinc.com</a>></span> wrote:<br>
<blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex"># HG changeset patch<br>
# User Wenju He <<a href="mailto:wenju@multicorewareinc.com">wenju@multicorewareinc.com</a>><br>
# Date 1379443287 18000<br>
# Tue Sep 17 13:41:27 2013 -0500<br>
# Node ID 97e1bed019a62849c5cc514ee74efbf64647986c<br>
# Parent 0d33ff236f68bc2238138a7213301b2efc0e6426<br>
add OpenCL environment codes<br>
<br>
diff -r 0d33ff236f68 -r 97e1bed019a6 source/CMakeLists.txt<br>
--- a/source/CMakeLists.txt Mon Sep 16 21:06:49 2013 +0800<br>
+++ b/source/CMakeLists.txt Tue Sep 17 13:41:27 2013 -0500<br>
@@ -1,190 +1,198 @@<br>
-if(NOT CMAKE_BUILD_TYPE)<br></blockquote><div><br></div><div>you've changed the file EOLN inadvertently, causing every line to be changed.</div><div> </div><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">
- # default to Release build for GCC builds<br>
- set(CMAKE_BUILD_TYPE Release CACHE STRING<br>
- "Choose the type of build, options are: None(CMAKE_CXX_FLAGS or CMAKE_C_FLAGS used) Debug Release RelWithDebInfo MinSizeRel."<br>
- FORCE)<br>
-endif()<br>
-<br>
-project (x265)<br>
-cmake_minimum_required (VERSION 2.6)<br>
-<br>
-SET(CMAKE_MODULE_PATH "${PROJECT_SOURCE_DIR}/cmake" "${CMAKE_MODULE_PATH}")<br>
-<br>
-if ("${CMAKE_SIZEOF_VOID_P}" MATCHES 8)<br>
- set(X64 1)<br>
- add_definitions(-DX86_64=1)<br>
-endif()<br>
-<br>
-if (CMAKE_GENERATOR STREQUAL "Xcode")<br>
- set(XCODE 1)<br>
-endif()<br>
-if (APPLE)<br>
- add_definitions(-DMACOS)<br>
-endif()<br>
-<br>
-# Enforce coding standards. Full warnings and warnings as errors<br>
-if("$ENV{CXX}" STREQUAL "icl")<br>
- set(MSVC 1)<br>
-endif()<br>
-if(MSVC)<br>
- option(STATIC_LINK_CRT "Statically link C runtime for release builds" OFF)<br>
- if (STATIC_LINK_CRT)<br>
- set(CompilerFlags CMAKE_CXX_FLAGS_RELEASE CMAKE_C_FLAGS_RELEASE)<br>
- foreach(CompilerFlag ${CompilerFlags})<br>
- string(REPLACE "/MD" "/MT" ${CompilerFlag} "${${CompilerFlag}}")<br>
- endforeach()<br>
- endif (STATIC_LINK_CRT)<br>
- add_definitions(/W4 /WX /D_CRT_SECURE_NO_WARNINGS)<br>
- add_definitions(/Ob2) # always inline<br>
- add_definitions(/Oi) # enable intrinsics<br>
- add_definitions(/MP) # multithreaded build<br>
- include_directories(compat/msvc)<br>
-endif(MSVC)<br>
-<br>
-if("$ENV{CXX}" STREQUAL "icpc")<br>
- set(GCC 1)<br>
- add_definitions(-Wall -Wextra -Wshadow -no-vec)<br>
-elseif(CMAKE_COMPILER_IS_GNUCXX)<br>
- execute_process(COMMAND ${CMAKE_CXX_COMPILER} -dumpversion OUTPUT_VARIABLE GCC_VERSION)<br>
- set(GCC 1)<br>
- add_definitions(-Wall -Wextra -Wshadow -mstackrealign)<br>
- if(NOT GCC_VERSION VERSION_LESS 4.7)<br>
- # this is necessary to avoid name conflicts in vector class<br>
- # library. if vector classes are removed/replaced this can<br>
- # likely be removed as well.<br>
- add_definitions(-fabi-version=6)<br>
- endif()<br>
-elseif(APPLE)<br>
- exec_program(uname ARGS -v OUTPUT_VARIABLE DARWIN_VERSION)<br>
- string(REGEX MATCH "[0-9]+" DARWIN_VERSION ${DARWIN_VERSION})<br>
- message(STATUS "DARWIN_VERSION=${DARWIN_VERSION}")<br>
- if (DARWIN_VERSION GREATER 12)<br>
- # Detect Mac OS X 10.8 llvm/gcc frankenstein's monster<br>
- if (${CMAKE_CXX_COMPILER} STREQUAL "/usr/bin/c++")<br>
- execute_process(COMMAND ${CMAKE_CXX_COMPILER} -dumpversion OUTPUT_VARIABLE GCC_VERSION)<br>
- set(GCC 1)<br>
- add_definitions(-Wall -Wextra -Wshadow)<br>
- endif()<br>
- endif()<br>
-endif()<br>
-if (GCC)<br>
- option(WARNINGS_AS_ERRORS "Stop compiles on first warning" OFF)<br>
- if(WARNINGS_AS_ERRORS)<br>
- add_definitions(-Werror)<br>
- endif(WARNINGS_AS_ERRORS)<br>
- if(X64 AND NOT WIN32)<br>
- set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -fPIC")<br>
- endif(X64 AND NOT WIN32)<br>
-endif(GCC)<br>
-<br>
-option(HIGH_BIT_DEPTH "Use 16bit pixels internally" OFF)<br>
-if(HIGH_BIT_DEPTH)<br>
- add_definitions(-DHIGH_BIT_DEPTH=1)<br>
-else(HIGH_BIT_DEPTH)<br>
- add_definitions(-DHIGH_BIT_DEPTH=0)<br>
-endif(HIGH_BIT_DEPTH)<br>
-<br>
-option(ENABLE_PRIMITIVES_VEC "Enable use of SIMD vector class primitives" ON)<br>
-find_package(Yasm)<br>
-if(YASM_FOUND)<br>
- if (YASM_VERSION_STRING VERSION_LESS "1.2.0")<br>
- message(STATUS "Yasm version ${YASM_VERSION_STRING} is too old. 1.2.0 or later required")<br>
- else()<br>
- message(STATUS "Found Yasm ${YASM_VERSION_STRING} to build assembly primitives")<br>
- option(ENABLE_PRIMITIVES_ASM "Enable use of assembly coded primitives" ON)<br>
- endif()<br>
-endif(YASM_FOUND)<br>
-<br>
-if(UNIX)<br>
- SET(PLATFORM_LIBS pthread m)<br>
- if(NOT ${CMAKE_SYSTEM_NAME} MATCHES "Darwin")<br>
- SET(PLATFORM_LIBS ${PLATFORM_LIBS} rt)<br>
- endif(NOT ${CMAKE_SYSTEM_NAME} MATCHES "Darwin")<br>
-endif(UNIX)<br>
-<br>
-option(ENABLE_STLPORT "Enable use of STLPort STL replacement library" OFF)<br>
-if(ENABLE_STLPORT)<br>
- set(STLPORT_LOCATION "${STLPORT_INCLUDE_DIR}" CACHE PATH "Location of STLPort sources")<br>
- # STLPort headers must be included before system headers<br>
- include_directories(BEFORE ${STLPORT_LOCATION}/stlport)<br>
- link_directories(${STLPORT_LOCATION}/lib)<br>
- if("$ENV{CXX}" STREQUAL "icl")<br>
- add_definitions(/Qwd1879)<br>
- endif()<br>
-endif(ENABLE_STLPORT)<br>
-<br>
-option(ENABLE_PPA "Enable PPA profiling instrumentation" OFF)<br>
-if(ENABLE_PPA)<br>
- add_definitions(-DENABLE_PPA)<br>
- add_subdirectory(PPA)<br>
- SET(PLATFORM_LIBS ${PLATFORM_LIBS} PPA)<br>
- if(UNIX)<br>
- SET(PLATFORM_LIBS ${PLATFORM_LIBS} dl)<br>
- endif(UNIX)<br>
-endif(ENABLE_PPA)<br>
-<br>
-if (WIN32)<br>
- # Visual leak detector<br>
- find_package(VLD)<br>
- if(VLD_FOUND)<br>
- add_definitions(-DHAVE_VLD)<br>
- include_directories(${VLD_INCLUDE_DIRS})<br>
- set(PLATFORM_LIBS ${PLATFORM_LIBS} ${VLD_LIBRARIES})<br>
- link_directories(${VLD_LIBRARY_DIRS})<br>
- endif()<br>
- option(WINXP_SUPPORT "Make binaries compatible with Windows XP" OFF)<br>
- if(WINXP_SUPPORT)<br>
- # force workarounds for atomic intrinsics introduced after XP<br>
- add_definitions(-D_WIN32_WINNT=_WIN32_WINNT_WINXP)<br>
- endif(WINXP_SUPPORT)<br>
-endif()<br>
-<br>
-include_directories(. Lib common encoder)<br>
-add_subdirectory(common)<br>
-add_subdirectory(encoder)<br>
-<br>
-set(LIBS common encoder)<br>
-if(ENABLE_PRIMITIVES_VEC)<br>
- set(LIBS ${LIBS} PrimitivesVec)<br>
-endif(ENABLE_PRIMITIVES_VEC)<br>
-<br>
-if(ENABLE_PRIMITIVES_ASM)<br>
- set(LIBS ${LIBS} PrimitivesASM)<br>
-endif(ENABLE_PRIMITIVES_ASM)<br>
-<br>
-if(NOT XCODE)<br>
- include(mergestaticlibs)<br>
- merge_static_libs(x265 ${LIBS})<br>
-endif()<br>
-<br>
-# Test applications<br>
-option(ENABLE_TESTS "Enable Unit Tests" OFF)<br>
-if(ENABLE_TESTS)<br>
- add_subdirectory(test)<br>
-endif(ENABLE_TESTS)<br>
-<br>
-# Main CLI application<br>
-option(ENABLE_CLI "Build standalone CLI application" ON)<br>
-if(ENABLE_CLI)<br>
- file(GLOB InputFiles input/*.cpp input/*.h)<br>
- file(GLOB OutputFiles output/*.cpp output/*.h)<br>
- source_group(input FILES ${InputFiles})<br>
- source_group(output FILES ${OutputFiles})<br>
- if (MSVC)<br>
- set_source_files_properties(compat/msvc/getopt.c PROPERTIES COMPILE_FLAGS "/wd4100")<br>
- elseif(GCC)<br>
- set_source_files_properties(compat/msvc/getopt.c PROPERTIES COMPILE_FLAGS "-Wno-unused-parameter")<br>
- endif()<br>
-<br>
- include(version) # determine X265_VERSION<br>
- set_source_files_properties(x265.cpp PROPERTIES COMPILE_FLAGS -DX265_VERSION=${X265_VERSION})<br>
- add_executable(cli ${EXTRAS} ../COPYING ${InputFiles} ${OutputFiles}<br>
- x265.cpp x265opts.h x265.h<br>
- compat/msvc/getopt.c compat/msvc/getopt.h)<br>
-if(XCODE)<br>
- target_link_libraries(cli ${LIBS} ${PLATFORM_LIBS})<br>
-else()<br>
- target_link_libraries(cli x265 ${PLATFORM_LIBS})<br>
-endif()<br>
- SET_TARGET_PROPERTIES(cli PROPERTIES OUTPUT_NAME x265)<br>
-endif(ENABLE_CLI)<br>
+if(NOT CMAKE_BUILD_TYPE)<br>
+ # default to Release build for GCC builds<br>
+ set(CMAKE_BUILD_TYPE Release CACHE STRING<br>
+ "Choose the type of build, options are: None(CMAKE_CXX_FLAGS or CMAKE_C_FLAGS used) Debug Release RelWithDebInfo MinSizeRel."<br>
+ FORCE)<br>
+endif()<br>
+<br>
+project (x265)<br>
+cmake_minimum_required (VERSION 2.6)<br>
+<br>
+SET(CMAKE_MODULE_PATH "${PROJECT_SOURCE_DIR}/cmake" "${CMAKE_MODULE_PATH}")<br>
+<br>
+if ("${CMAKE_SIZEOF_VOID_P}" MATCHES 8)<br>
+ set(X64 1)<br>
+ add_definitions(-DX86_64=1)<br>
+endif()<br>
+<br>
+if (CMAKE_GENERATOR STREQUAL "Xcode")<br>
+ set(XCODE 1)<br>
+endif()<br>
+if (APPLE)<br>
+ add_definitions(-DMACOS)<br>
+endif()<br>
+<br>
+# Enforce coding standards. Full warnings and warnings as errors<br>
+if("$ENV{CXX}" STREQUAL "icl")<br>
+ set(MSVC 1)<br>
+endif()<br>
+if(MSVC)<br>
+ option(STATIC_LINK_CRT "Statically link C runtime for release builds" OFF)<br>
+ if (STATIC_LINK_CRT)<br>
+ set(CompilerFlags CMAKE_CXX_FLAGS_RELEASE CMAKE_C_FLAGS_RELEASE)<br>
+ foreach(CompilerFlag ${CompilerFlags})<br>
+ string(REPLACE "/MD" "/MT" ${CompilerFlag} "${${CompilerFlag}}")<br>
+ endforeach()<br>
+ endif (STATIC_LINK_CRT)<br>
+ add_definitions(/W4 /WX /D_CRT_SECURE_NO_WARNINGS)<br>
+ add_definitions(/Ob2) # always inline<br>
+ add_definitions(/Oi) # enable intrinsics<br>
+ add_definitions(/MP) # multithreaded build<br>
+ include_directories(compat/msvc)<br>
+endif(MSVC)<br>
+<br>
+if("$ENV{CXX}" STREQUAL "icpc")<br>
+ set(GCC 1)<br>
+ add_definitions(-Wall -Wextra -Wshadow -no-vec)<br>
+elseif(CMAKE_COMPILER_IS_GNUCXX)<br>
+ execute_process(COMMAND ${CMAKE_CXX_COMPILER} -dumpversion OUTPUT_VARIABLE GCC_VERSION)<br>
+ set(GCC 1)<br>
+ add_definitions(-Wall -Wextra -Wshadow -mstackrealign)<br>
+ if(NOT GCC_VERSION VERSION_LESS 4.7)<br>
+ # this is necessary to avoid name conflicts in vector class<br>
+ # library. if vector classes are removed/replaced this can<br>
+ # likely be removed as well.<br>
+ add_definitions(-fabi-version=6)<br>
+ endif()<br>
+elseif(APPLE)<br>
+ exec_program(uname ARGS -v OUTPUT_VARIABLE DARWIN_VERSION)<br>
+ string(REGEX MATCH "[0-9]+" DARWIN_VERSION ${DARWIN_VERSION})<br>
+ message(STATUS "DARWIN_VERSION=${DARWIN_VERSION}")<br>
+ if (DARWIN_VERSION GREATER 12)<br>
+ # Detect Mac OS X 10.8 llvm/gcc frankenstein's monster<br>
+ if (${CMAKE_CXX_COMPILER} STREQUAL "/usr/bin/c++")<br>
+ execute_process(COMMAND ${CMAKE_CXX_COMPILER} -dumpversion OUTPUT_VARIABLE GCC_VERSION)<br>
+ set(GCC 1)<br>
+ add_definitions(-Wall -Wextra -Wshadow)<br>
+ endif()<br>
+ endif()<br>
+endif()<br>
+if (GCC)<br>
+ option(WARNINGS_AS_ERRORS "Stop compiles on first warning" OFF)<br>
+ if(WARNINGS_AS_ERRORS)<br>
+ add_definitions(-Werror)<br>
+ endif(WARNINGS_AS_ERRORS)<br>
+ if(X64 AND NOT WIN32)<br>
+ set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -fPIC")<br>
+ endif(X64 AND NOT WIN32)<br>
+endif(GCC)<br>
+<br>
+option(HIGH_BIT_DEPTH "Use 16bit pixels internally" OFF)<br>
+if(HIGH_BIT_DEPTH)<br>
+ add_definitions(-DHIGH_BIT_DEPTH=1)<br>
+else(HIGH_BIT_DEPTH)<br>
+ add_definitions(-DHIGH_BIT_DEPTH=0)<br>
+endif(HIGH_BIT_DEPTH)<br>
+<br>
+option(ENABLE_PRIMITIVES_VEC "Enable use of SIMD vector class primitives" ON)<br>
+find_package(Yasm)<br>
+if(YASM_FOUND)<br>
+ if (YASM_VERSION_STRING VERSION_LESS "1.2.0")<br>
+ message(STATUS "Yasm version ${YASM_VERSION_STRING} is too old. 1.2.0 or later required")<br>
+ else()<br>
+ message(STATUS "Found Yasm ${YASM_VERSION_STRING} to build assembly primitives")<br>
+ option(ENABLE_PRIMITIVES_ASM "Enable use of assembly coded primitives" ON)<br>
+ endif()<br>
+endif(YASM_FOUND)<br>
+<br>
+if(UNIX)<br>
+ SET(PLATFORM_LIBS pthread m)<br>
+ if(NOT ${CMAKE_SYSTEM_NAME} MATCHES "Darwin")<br>
+ SET(PLATFORM_LIBS ${PLATFORM_LIBS} rt)<br>
+ endif(NOT ${CMAKE_SYSTEM_NAME} MATCHES "Darwin")<br>
+endif(UNIX)<br>
+<br>
+option(ENABLE_STLPORT "Enable use of STLPort STL replacement library" OFF)<br>
+if(ENABLE_STLPORT)<br>
+ set(STLPORT_LOCATION "${STLPORT_INCLUDE_DIR}" CACHE PATH "Location of STLPort sources")<br>
+ # STLPort headers must be included before system headers<br>
+ include_directories(BEFORE ${STLPORT_LOCATION}/stlport)<br>
+ link_directories(${STLPORT_LOCATION}/lib)<br>
+ if("$ENV{CXX}" STREQUAL "icl")<br>
+ add_definitions(/Qwd1879)<br>
+ endif()<br>
+endif(ENABLE_STLPORT)<br>
+<br>
+option(ENABLE_PPA "Enable PPA profiling instrumentation" OFF)<br>
+if(ENABLE_PPA)<br>
+ add_definitions(-DENABLE_PPA)<br>
+ add_subdirectory(PPA)<br>
+ SET(PLATFORM_LIBS ${PLATFORM_LIBS} PPA)<br>
+ if(UNIX)<br>
+ SET(PLATFORM_LIBS ${PLATFORM_LIBS} dl)<br>
+ endif(UNIX)<br>
+endif(ENABLE_PPA)<br>
+<br>
+option(ENABLE_OPENCL "Use OpenCL for motion estimation" OFF)<br>
+if(ENABLE_OPENCL)<br>
+ find_package(OpenCL)<br>
+ if(OPENCL_FOUND)<br>
+ add_definitions(-DHAVE_OPENCL)<br>
+ endif()<br>
+endif(ENABLE_OPENCL)<br>
+<br>
+if (WIN32)<br>
+ # Visual leak detector<br>
+ find_package(VLD)<br>
+ if(VLD_FOUND)<br>
+ add_definitions(-DHAVE_VLD)<br>
+ include_directories(${VLD_INCLUDE_DIRS})<br>
+ set(PLATFORM_LIBS ${PLATFORM_LIBS} ${VLD_LIBRARIES})<br>
+ link_directories(${VLD_LIBRARY_DIRS})<br>
+ endif()<br>
+ option(WINXP_SUPPORT "Make binaries compatible with Windows XP" OFF)<br>
+ if(WINXP_SUPPORT)<br>
+ # force workarounds for atomic intrinsics introduced after XP<br>
+ add_definitions(-D_WIN32_WINNT=_WIN32_WINNT_WINXP)<br>
+ endif(WINXP_SUPPORT)<br>
+endif()<br>
+<br>
+include_directories(. Lib common encoder)<br>
+add_subdirectory(common)<br>
+add_subdirectory(encoder)<br>
+<br>
+set(LIBS common encoder)<br>
+if(ENABLE_PRIMITIVES_VEC)<br>
+ set(LIBS ${LIBS} PrimitivesVec)<br>
+endif(ENABLE_PRIMITIVES_VEC)<br>
+<br>
+if(ENABLE_PRIMITIVES_ASM)<br>
+ set(LIBS ${LIBS} PrimitivesASM)<br>
+endif(ENABLE_PRIMITIVES_ASM)<br>
+<br>
+if(NOT XCODE)<br>
+ include(mergestaticlibs)<br>
+ merge_static_libs(x265 ${LIBS})<br>
+endif()<br>
+<br>
+# Test applications<br>
+option(ENABLE_TESTS "Enable Unit Tests" OFF)<br>
+if(ENABLE_TESTS)<br>
+ add_subdirectory(test)<br>
+endif(ENABLE_TESTS)<br>
+<br>
+# Main CLI application<br>
+option(ENABLE_CLI "Build standalone CLI application" ON)<br>
+if(ENABLE_CLI)<br>
+ file(GLOB InputFiles input/*.cpp input/*.h)<br>
+ file(GLOB OutputFiles output/*.cpp output/*.h)<br>
+ source_group(input FILES ${InputFiles})<br>
+ source_group(output FILES ${OutputFiles})<br>
+ if (MSVC)<br>
+ set_source_files_properties(compat/msvc/getopt.c PROPERTIES COMPILE_FLAGS "/wd4100")<br>
+ elseif(GCC)<br>
+ set_source_files_properties(compat/msvc/getopt.c PROPERTIES COMPILE_FLAGS "-Wno-unused-parameter")<br>
+ endif()<br>
+<br>
+ include(version) # determine X265_VERSION<br>
+ set_source_files_properties(x265.cpp PROPERTIES COMPILE_FLAGS -DX265_VERSION=${X265_VERSION})<br>
+ add_executable(cli ${EXTRAS} ../COPYING ${InputFiles} ${OutputFiles}<br>
+ x265.cpp x265opts.h x265.h<br>
+ compat/msvc/getopt.c compat/msvc/getopt.h)<br>
+if(XCODE)<br>
+ target_link_libraries(cli ${LIBS} ${PLATFORM_LIBS} ${OPENCL_LIBRARIES})<br>
+else()<br>
+ target_link_libraries(cli x265 ${PLATFORM_LIBS} ${OPENCL_LIBRARIES})<br>
+endif()<br>
+ SET_TARGET_PROPERTIES(cli PROPERTIES OUTPUT_NAME x265)<br>
+endif(ENABLE_CLI)<br>
diff -r 0d33ff236f68 -r 97e1bed019a6 source/cmake/FindOpenCL.cmake<br>
--- /dev/null Thu Jan 01 00:00:00 1970 +0000<br>
+++ b/source/cmake/FindOpenCL.cmake Tue Sep 17 13:41:27 2013 -0500<br>
@@ -0,0 +1,58 @@<br>
+# modified from OpenCV<br></blockquote><div><br></div><div>If this file was copied and adapted from OpenCV, we need to keep their license header and attribution here, even if MCW contributed the patch.</div><div><br></div>
<blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">
+<br>
+if(APPLE)<br>
+ set(OPENCL_FOUND YES)<br>
+ set(OPENCL_LIBRARY "-framework OpenCL" CACHE STRING "OpenCL library")<br>
+ set(OPENCL_INCLUDE_DIR "" CACHE STRING "OpenCL include directory")<br>
+ mark_as_advanced(OPENCL_INCLUDE_DIR OPENCL_LIBRARY)<br>
+else(APPLE)<br>
+ #find_package(OpenCL QUIET)<br></blockquote><div><br></div><div>let's not add commented lines in new files</div><div> </div><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">
+<br>
+ if (NOT OPENCL_FOUND)<br>
+ find_path(OPENCL_ROOT_DIR<br>
+ NAMES OpenCL/cl.h CL/cl.h include/CL/cl.h include/nvidia-current/CL/cl.h<br>
+ PATHS ENV OCLROOT ENV AMDAPPSDKROOT ENV CUDA_PATH ENV INTELOCLSDKROOT<br>
+ DOC "OpenCL root directory"<br>
+ NO_DEFAULT_PATH)<br>
+<br>
+ find_path(OPENCL_INCLUDE_DIR<br>
+ NAMES OpenCL/cl.h CL/cl.h<br>
+ HINTS ${OPENCL_ROOT_DIR}<br>
+ PATH_SUFFIXES include include/nvidia-current<br>
+ DOC "OpenCL include directory"<br>
+ NO_DEFAULT_PATH)<br>
+<br>
+ if(WIN32)<br></blockquote><div><br></div><div>our cmake script style uses 4-spaces for intendation</div><div> </div><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">
+ if(X64)<br>
+ set(OPENCL_POSSIBLE_LIB_SUFFIXES lib/Win64 lib/x86_64 lib/x64)<br>
+ else()<br>
+ set(OPENCL_POSSIBLE_LIB_SUFFIXES lib/Win32 lib/x86)<br>
+ endif()<br>
+ elseif(UNIX)<br>
+ if(X64)<br>
+ set(OPENCL_POSSIBLE_LIB_SUFFIXES lib64 lib)<br>
+ else()<br>
+ set(OPENCL_POSSIBLE_LIB_SUFFIXES lib32 lib)<br>
+ endif()<br>
+ else()<br>
+ set(OPENCL_POSSIBLE_LIB_SUFFIXES lib)<br>
+ endif()<br>
+<br>
+ find_library(OPENCL_LIBRARY<br>
+ NAMES OpenCL<br>
+ HINTS ${OPENCL_ROOT_DIR}<br>
+ PATH_SUFFIXES ${OPENCL_POSSIBLE_LIB_SUFFIXES}<br>
+ DOC "OpenCL library"<br>
+ NO_DEFAULT_PATH)<br>
+<br>
+ mark_as_advanced(OPENCL_INCLUDE_DIR OPENCL_LIBRARY)<br>
+ include(FindPackageHandleStandardArgs)<br>
+ FIND_PACKAGE_HANDLE_STANDARD_ARGS(OPENCL DEFAULT_MSG OPENCL_LIBRARY OPENCL_INCLUDE_DIR )<br>
+ endif()<br>
+endif(APPLE)<br>
+<br>
+if(OPENCL_FOUND)<br>
+ set(HAVE_OPENCL 1)<br>
+ set(OPENCL_INCLUDE_DIRS ${OPENCL_INCLUDE_DIR})<br>
+ set(OPENCL_LIBRARIES ${OPENCL_LIBRARY})<br>
+endif()<br>
diff -r 0d33ff236f68 -r 97e1bed019a6 source/cmake/cl2cpp.cmake<br>
--- /dev/null Thu Jan 01 00:00:00 1970 +0000<br>
+++ b/source/cmake/cl2cpp.cmake Tue Sep 17 13:41:27 2013 -0500<br>
@@ -0,0 +1,35 @@<br>
+# modified from OpenCV<br></blockquote><div><br></div><div>ditto attribution here</div><div> </div><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">
+<br>
+file(GLOB cl_list "${CL_DIR}/*.h" "${CL_DIR}/*.cl" )<br>
+<br>
+file(WRITE ${OUTPUT} "// This file is auto-generated. Do not edit!<br>
+<br>
+namespace x265<br>
+{<br>
+<br>
+")<br>
+<br>
+file(APPEND ${OUTPUT} "const char* x265_opencl_source=\n")<br>
+<br>
+foreach(cl ${cl_list})<br>
+ file(READ "${cl}" lines)<br>
+<br>
+ string(REPLACE "\r" "" lines "${lines}\n")<br>
+ string(REPLACE "\t" " " lines "${lines}")<br>
+<br>
+ string(REGEX REPLACE "/\\*([^*]/|\\*[^/]|[^*/])*\\*/" "" lines "${lines}") # multiline comments<br>
+ string(REGEX REPLACE "/\\*([^\n])*\\*/" "" lines "${lines}") # single-line comments<br>
+ string(REGEX REPLACE "[ ]*//[^\n]*\n" "\n" lines "${lines}") # single-line comments<br>
+ string(REGEX REPLACE "\n[ ]*(\n[ ]*)*" "\n" lines "${lines}") # empty lines & leading whitespace<br>
+ string(REGEX REPLACE "^\n" "" lines "${lines}") # leading new line<br>
+<br>
+ string(REPLACE "\\" "\\\\" lines "${lines}")<br>
+ string(REPLACE "\"" "\\\"" lines "${lines}")<br>
+ string(REPLACE "\n" "\\n\"\n\"" lines "${lines}")<br>
+<br>
+ string(REGEX REPLACE "\"$" "" lines "${lines}") # unneeded " at the eof<br>
+<br>
+ file(APPEND ${OUTPUT} "\"${lines}")<br>
+endforeach()<br>
+<br>
+file(APPEND ${OUTPUT} ";\n\n}\n")<br></blockquote><div><br></div><div>we should also cleanup end of file eoln</div><div> </div><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">
\ No newline at end of file<br>
diff -r 0d33ff236f68 -r 97e1bed019a6 source/common/CMakeLists.txt<br>
--- a/source/common/CMakeLists.txt Mon Sep 16 21:06:49 2013 +0800<br>
+++ b/source/common/CMakeLists.txt Tue Sep 17 13:41:27 2013 -0500<br>
@@ -1,59 +1,73 @@<br>
-if(GCC)<br></blockquote><div><br></div><div>again, you've changed the EOLN of this file on accident</div><div> </div><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">
- if (NOT X64)<br>
- # force gcc to generate code for sync primitives<br>
- set_source_files_properties(threadpool.cpp reference.cpp wavefront.cpp PROPERTIES COMPILE_FLAGS -march=i686)<br>
- endif()<br>
-endif(GCC)<br>
-<br>
-if(ENABLE_PRIMITIVES_VEC)<br>
- add_definitions(-DENABLE_VECTOR_PRIMITIVES=1)<br>
-endif(ENABLE_PRIMITIVES_VEC)<br>
-if(ENABLE_PRIMITIVES_ASM)<br>
- add_definitions(-DENABLE_ASM_PRIMITIVES=1)<br>
-endif(ENABLE_PRIMITIVES_ASM)<br>
-<br>
-include_directories(../VectorClass)<br>
-<br>
-file(GLOB LIBCOMMON_HDR ../Lib/TLibCommon/*.h)<br>
-file(GLOB LIBCOMMON_SRC ../Lib/TLibCommon/*.cpp)<br>
-source_group(TLibCommon FILES ${LIBCOMMON_SRC})<br>
-source_group(TLibCommonH FILES ${LIBCOMMON_HDR})<br>
-if(GCC)<br>
- set_source_files_properties(${LIBCOMMON_SRC} PROPERTIES COMPILE_FLAGS<br>
- "-Wno-sign-compare")<br>
-endif(GCC)<br>
-if(MSVC)<br>
- # ignore these warnings from HM source<br>
- # /wd4244 type conversion, possible loss of data<br>
- # /wd4512 assignment operator could not be generated<br>
- # /wd4127 conditional expression is constant<br>
- # /wd4389 signed/unsigned mismatch<br>
- # /wd4018 '<' signed/unsigned mismatch<br>
- # /wd4800 performance warning: bool coersion<br>
- set_source_files_properties(${LIBCOMMON_SRC} PROPERTIES COMPILE_FLAGS<br>
- "/wd4244 /wd4512 /wd4127 /wd4389 /wd4018 /wd4800")<br>
- if ("$ENV{CXX}" STREQUAL "icl")<br>
- add_definitions(/Qwd2557) # signed/unsigned mismatch<br>
- endif()<br>
-endif(MSVC)<br>
-<br>
-add_library(common STATIC ../../COPYING<br>
- ${LIBCOMMON_SRC} ${LIBCOMMON_HDR}<br>
- primitives.cpp primitives.h<br>
- pixel.cpp dct.cpp ipfilter.cpp intrapred.cpp<br>
- ../VectorClass/instrset_detect.cpp<br>
- threading.cpp threading.h<br>
- threadpool.cpp threadpool.h<br>
- wavefront.h wavefront.cpp<br>
- md5.cpp md5.h<br>
- TShortYUV.cpp TShortYUV.h mv.h<br>
- reference.cpp reference.h<br>
- common.cpp common.h<br>
- lowres.cpp lowres.h)<br>
-<br>
-if(ENABLE_PRIMITIVES_VEC)<br>
- add_subdirectory(vec)<br>
-endif(ENABLE_PRIMITIVES_VEC)<br>
-if(ENABLE_PRIMITIVES_ASM)<br>
- add_subdirectory(x86)<br>
-endif(ENABLE_PRIMITIVES_ASM)<br>
+if(GCC)<br>
+ if (NOT X64)<br>
+ # force gcc to generate code for sync primitives<br>
+ set_source_files_properties(threadpool.cpp reference.cpp wavefront.cpp PROPERTIES COMPILE_FLAGS -march=i686)<br>
+ endif()<br>
+endif(GCC)<br>
+<br>
+if(ENABLE_PRIMITIVES_VEC)<br>
+ add_definitions(-DENABLE_VECTOR_PRIMITIVES=1)<br>
+endif(ENABLE_PRIMITIVES_VEC)<br>
+if(ENABLE_PRIMITIVES_ASM)<br>
+ add_definitions(-DENABLE_ASM_PRIMITIVES=1)<br>
+endif(ENABLE_PRIMITIVES_ASM)<br>
+<br>
+include_directories(../VectorClass)<br>
+<br>
+file(GLOB LIBCOMMON_HDR ../Lib/TLibCommon/*.h)<br>
+file(GLOB LIBCOMMON_SRC ../Lib/TLibCommon/*.cpp)<br>
+source_group(TLibCommon FILES ${LIBCOMMON_SRC})<br>
+source_group(TLibCommonH FILES ${LIBCOMMON_HDR})<br>
+if(GCC)<br>
+ set_source_files_properties(${LIBCOMMON_SRC} PROPERTIES COMPILE_FLAGS<br>
+ "-Wno-sign-compare")<br>
+endif(GCC)<br>
+if(MSVC)<br>
+ # ignore these warnings from HM source<br>
+ # /wd4244 type conversion, possible loss of data<br>
+ # /wd4512 assignment operator could not be generated<br>
+ # /wd4127 conditional expression is constant<br>
+ # /wd4389 signed/unsigned mismatch<br>
+ # /wd4018 '<' signed/unsigned mismatch<br>
+ # /wd4800 performance warning: bool coersion<br>
+ set_source_files_properties(${LIBCOMMON_SRC} PROPERTIES COMPILE_FLAGS<br>
+ "/wd4244 /wd4512 /wd4127 /wd4389 /wd4018 /wd4800")<br>
+ if ("$ENV{CXX}" STREQUAL "icl")<br>
+ add_definitions(/Qwd2557) # signed/unsigned mismatch<br>
+ endif()<br>
+endif(MSVC)<br>
+<br>
+if(HAVE_OPENCL)<br>
+ include_directories(${OPENCL_INCLUDE_DIRS})<br>
+ set(OPENCL_KERNEL_HDR ${CMAKE_CURRENT_SOURCE_DIR}/oclkernels.h)<br>
+ add_custom_command(<br>
+ OUTPUT ${OPENCL_KERNEL_HDR}<br>
+ COMMAND ${CMAKE_COMMAND} -DCL_DIR="${CMAKE_CURRENT_SOURCE_DIR}/opencl"<br>
+ -DOUTPUT=${OPENCL_KERNEL_HDR}<br>
+ -P "${CMAKE_CURRENT_SOURCE_DIR}/../cmake/cl2cpp.cmake"<br>
+ DEPENDS "${CMAKE_CURRENT_SOURCE_DIR}/../cmake/cl2cpp.cmake")<br>
+ set(OPENCL_SRC opencl.cpp)<br>
+ set(OPENCL_HDR opencl.h ${OPENCL_KERNEL_HDR})<br>
+endif(HAVE_OPENCL)<br>
+<br>
+add_library(common STATIC ../../COPYING<br>
+ ${LIBCOMMON_SRC} ${LIBCOMMON_HDR}<br>
+ primitives.cpp primitives.h<br>
+ pixel.cpp dct.cpp ipfilter.cpp intrapred.cpp<br>
+ ../VectorClass/instrset_detect.cpp<br>
+ threading.cpp threading.h<br>
+ threadpool.cpp threadpool.h<br>
+ wavefront.h wavefront.cpp<br>
+ md5.cpp md5.h<br>
+ TShortYUV.cpp TShortYUV.h mv.h<br>
+ reference.cpp reference.h<br>
+ common.cpp common.h<br>
+ lowres.cpp lowres.h<br>
+ ${OPENCL_SRC} ${OPENCL_HDR})<br>
+<br>
+if(ENABLE_PRIMITIVES_VEC)<br>
+ add_subdirectory(vec)<br>
+endif(ENABLE_PRIMITIVES_VEC)<br>
+if(ENABLE_PRIMITIVES_ASM)<br>
+ add_subdirectory(x86)<br>
+endif(ENABLE_PRIMITIVES_ASM)<br>
diff -r 0d33ff236f68 -r 97e1bed019a6 source/common/opencl.cpp<br>
--- /dev/null Thu Jan 01 00:00:00 1970 +0000<br>
+++ b/source/common/opencl.cpp Tue Sep 17 13:41:27 2013 -0500<br>
@@ -0,0 +1,631 @@<br>
+/*****************************************************************************<br>
+ * opencl.c: OpenCL initialization and kernel compilation<br>
+ *****************************************************************************<br>
+ * Copyright (C) 2013 x265 project<br>
+ *<br>
+ * Authors: Steve Borho <<a href="mailto:steve@borho.org">steve@borho.org</a>><br>
+ *<br>
+ * This program is free software; you can redistribute it and/or modify<br>
+ * it under the terms of the GNU General Public License as published by<br>
+ * the Free Software Foundation; either version 2 of the License, or<br>
+ * (at your option) any later version.<br>
+ *<br>
+ * This program is distributed in the hope that it will be useful,<br>
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of<br>
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the<br>
+ * GNU General Public License for more details.<br>
+ *<br>
+ * You should have received a copy of the GNU General Public License<br>
+ * along with this program; if not, write to the Free Software<br>
+ * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02111, USA.<br>
+ *<br>
+ * This program is also available under a commercial proprietary license.<br>
+ * For more information, contact us at <a href="mailto:licensing@multicorewareinc.com">licensing@multicorewareinc.com</a>.<br>
+ *****************************************************************************/<br>
+<br>
+#include "TLibCommon/CommonDef.h"<br>
+#include "common.h"<br>
+<br>
+#if defined WIN32 || defined _WIN32<br>
+#include <Windows.h><br></blockquote><div><br></div><div>windows.h needs to be lower case, this actually matters for mingw cross compiles</div><div> </div><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">
+#else<br>
+#include <dlfcn.h> //dlopen, dlsym, dlclose<br>
+#endif<br>
+#include "opencl.h"<br>
+#include "oclkernels.h"<br>
+#include "stdio.h"<br>
+<br>
+/* define from recent cl_ext.h, copied here in case headers are old */<br>
+#define CL_DEVICE_SIMD_INSTRUCTION_WIDTH_AMD 0x4042<br>
+<br>
+using namespace x265;<br>
+<br>
+// global variable<br>
+static ocl g_ocl;<br>
+<br>
+static void CALLBACK oclErrorNotify( const char *errinfo, const void *private_info, size_t cb, void *user_data )<br>
+{<br>
+ /* Any error notification can be assumed to be fatal to the OpenCL context.<br>
+ * We need to stop using it immediately to prevent further damage. */<br>
+ x265_param_t *param = (x265_param_t*)user_data;<br>
+ param->useOpenCL = 0;<br>
+ //h->opencl.b_fatal_error = 1;<br>
+ x265_log(param, X265_LOG_ERROR, "OpenCL: %s\n", errinfo );<br>
+ x265_log(param, X265_LOG_ERROR, "OpenCL: fatal error, aborting encode\n" );<br>
+ (void)cb;<br>
+ (void)private_info;<br>
+}<br>
+<br>
+ocl::ocl()<br>
+ : m_context(NULL)<br>
+ , m_device(NULL)<br>
+ , m_queue(NULL)<br>
+ , m_program(NULL)<br>
+ , m_initialized(false)<br>
+ , m_deviceAMDSI(false)<br>
+ , m_fatalError(false)<br>
+ , m_openclThreadPri(0)<br>
+ , m_kernelMeStar(NULL)<br>
+{}<br>
+<br>
+ocl::~ocl()<br>
+{}<br>
+<br>
+<br>
+bool ocl::init(x265_param_t *param)<br>
+{<br>
+ cl_int status;<br>
+ cl_uint numPlatforms;<br>
+ bool ret = false;<br>
+<br>
+ if (m_initialized)<br>
+ {<br>
+ return true;<br>
+ }<br>
+<br>
+ status = clGetPlatformIDs(0, NULL, &numPlatforms);<br>
+ if( status != CL_SUCCESS || numPlatforms == 0 )<br>
+ {<br>
+ x265_log(param, X265_LOG_WARNING, "OpenCL: Unable to query installed platforms\n");<br>
+ return false;<br>
+ }<br>
+<br>
+ cl_platform_id *platforms = (cl_platform_id*)X265_MALLOC(cl_platform_id, numPlatforms);<br>
+ status = clGetPlatformIDs( numPlatforms, platforms, NULL );<br>
+ if( status != CL_SUCCESS )<br>
+ {<br>
+ x265_log(param, X265_LOG_WARNING, "OpenCL: Unable to query installed platforms\n");<br>
+ X265_FREE( platforms );<br>
+ return false;<br>
+ }<br>
+<br>
+ /* Select the first OpenCL platform with a GPU device that supports our<br>
+ * required image (texture) formats */<br>
+ for( cl_uint i = 0; i < numPlatforms; ++i )<br>
+ {<br>
+ cl_uint gpu_count = 0;<br>
+ status = clGetDeviceIDs( platforms[i], CL_DEVICE_TYPE_GPU, 0, NULL, &gpu_count );<br>
+ if( status != CL_SUCCESS || !gpu_count )<br>
+ continue;<br>
+<br>
+ cl_device_id *devices = (cl_device_id*) X265_MALLOC( cl_device_id, gpu_count );<br>
+ if( !devices )<br>
+ continue;<br>
+<br>
+ status = clGetDeviceIDs( platforms[i], CL_DEVICE_TYPE_GPU, gpu_count, devices, NULL );<br>
+ if( status != CL_SUCCESS )<br>
+ {<br>
+ X265_FREE( devices );<br>
+ continue;<br>
+ }<br>
+<br>
+ /* Find a GPU device that supports our image formats */<br>
+ for( cl_uint gpu = 0; gpu < gpu_count; gpu++ )<br>
+ {<br>
+ cl_device_id device = devices[gpu];<br>
+<br>
+ /* if the user has specified an exact device ID, skip all other<br>
+ * GPUs. If this device matches, allow it to continue through the<br>
+ * checks for supported images, etc. */<br>
+ if( param->OpenCLDeviceId && devices[gpu] != (cl_device_id) param->OpenCLDeviceId )<br>
+ continue;<br>
+<br>
+ cl_bool image_support;<br>
+ clGetDeviceInfo( device, CL_DEVICE_IMAGE_SUPPORT, sizeof(cl_bool), &image_support, NULL );<br>
+ if( !image_support )<br>
+ continue;<br>
+<br>
+ cl_context context = clCreateContext( NULL, 1, &device, oclErrorNotify, (void*)param, &status );<br>
+ if( status != CL_SUCCESS )<br>
+ continue;<br>
+<br>
+ cl_uint imagecount = 0;<br>
+ clGetSupportedImageFormats( context, CL_MEM_READ_WRITE, CL_MEM_OBJECT_IMAGE2D, 0, NULL, &imagecount );<br>
+ if( !imagecount )<br>
+ {<br>
+ clReleaseContext( context );<br>
+ continue;<br>
+ }<br>
+<br>
+ cl_image_format *imageType = (cl_image_format*)X265_MALLOC(cl_image_format, imagecount);<br>
+ if( !imageType )<br>
+ {<br>
+ clReleaseContext( context );<br>
+ continue;<br>
+ }<br>
+<br>
+ clGetSupportedImageFormats( context, CL_MEM_READ_WRITE, CL_MEM_OBJECT_IMAGE2D, imagecount, imageType, NULL );<br>
+<br>
+ int b_has_r = 0;<br>
+ int b_has_rgba = 0;<br>
+ for( cl_uint j = 0; j < imagecount; j++ )<br>
+ {<br>
+ if( imageType[j].image_channel_order == CL_R &&<br>
+ imageType[j].image_channel_data_type == CL_UNSIGNED_INT32 )<br>
+ b_has_r = 1;<br>
+ else if( imageType[j].image_channel_order == CL_RGBA &&<br>
+ imageType[j].image_channel_data_type == CL_UNSIGNED_INT8 )<br>
+ b_has_rgba = 1;<br>
+ }<br>
+ X265_FREE( imageType );<br>
+ if( !b_has_r || !b_has_rgba )<br>
+ {<br>
+ char devname[64];<br>
+ status = clGetDeviceInfo(device, CL_DEVICE_NAME, sizeof(devname), devname, NULL);<br>
+ if( status == CL_SUCCESS )<br>
+ {<br>
+ /* emit warning if we are discarding the user's explicit choice */<br>
+ int level = param->OpenCLDeviceId ? X265_LOG_WARNING : X265_LOG_DEBUG;<br>
+ x265_log(param, level, "OpenCL: %s does not support required image formats\n", devname);<br>
+ }<br>
+ clReleaseContext( context );<br>
+ continue;<br>
+ }<br>
+<br>
+ /* user selection of GPU device, skip N first matches */<br>
+ if( param->iOpenCLDevice )<br>
+ {<br>
+ param->iOpenCLDevice--;<br>
+ clReleaseContext( context );<br>
+ continue;<br>
+ }<br>
+<br>
+ cl_command_queue queue = clCreateCommandQueue( context, device, 0, &status );<br>
+ if( status != CL_SUCCESS )<br></blockquote><div><br></div><div>watch the white-space coding style</div><div> </div><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">
+ {<br>
+ clReleaseContext( context );<br>
+ continue;<br>
+ }<br>
+<br>
+ m_context = context;<br>
+ m_device = device;<br>
+ m_queue = queue;<br>
+<br>
+ ret = true;<br>
+ break;<br>
+ }<br>
+<br>
+ X265_FREE( devices );<br>
+<br>
+ if (ret)<br>
+ {<br>
+ break;<br>
+ }<br>
+ }<br>
+<br>
+ X265_FREE( platforms );<br>
+<br>
+ if (!param->pszCLbinFile)<br>
+ {<br>
+ param->pszCLbinFile = "x265_kernels.clbin";<br>
+ }<br>
+<br>
+ if (!ret)<br>
+ {<br>
+ x265_log(param, X265_LOG_WARNING, "OpenCL: Unable to find a compatible device\n");<br>
+ }<br>
+ else<br>
+ {<br>
+ ret = createKernels(param);<br>
+ }<br>
+<br>
+ m_initialized = ret;<br>
+<br>
+ return ret;<br>
+}<br>
+<br>
+bool ocl::createKernels(x265_param_t *param)<br>
+{<br>
+ static const char *kernelnames[] = {<br>
+ "motion_estimate"<br>
+ };<br>
+ cl_kernel *kernels[] = {<br>
+ &m_kernelMeStar<br>
+ };<br>
+ cl_int status;<br>
+<br>
+ compile(param);<br>
+ if( !m_program )<br>
+ {<br>
+ return false;<br>
+ }<br>
+<br>
+ for( int i = 0; i < sizeof(kernelnames)/sizeof(kernelnames[0]); i++ )<br>
+ {<br>
+ *kernels[i] = clCreateKernel(m_program, kernelnames[i], &status);<br>
+ if( status != CL_SUCCESS )<br>
+ {<br>
+ x265_log(param, X265_LOG_WARNING, "OpenCL: Unable to compile kernel '%s' (%d)\n", kernelnames[i], status );<br>
+ return false;<br>
+ }<br>
+ }<br>
+<br>
+ return true;<br>
+}<br>
+<br>
+void ocl::destroy()<br>
+{<br>
+ if (m_queue)<br>
+ {<br>
+ clFinish(m_queue);<br>
+ }<br>
+<br>
+ if (m_kernelMeStar)<br>
+ {<br>
+ clReleaseKernel(m_kernelMeStar);<br>
+ m_kernelMeStar = NULL;<br>
+ }<br>
+<br>
+ if (m_program)<br>
+ {<br>
+ clReleaseProgram(m_program);<br>
+ m_program = NULL;<br>
+ }<br>
+<br>
+ if (m_queue)<br>
+ {<br>
+ clReleaseCommandQueue( m_queue );<br>
+ m_queue = NULL;<br>
+ }<br>
+<br>
+ if (m_context)<br>
+ {<br>
+ clReleaseContext( m_context );<br>
+ m_context = NULL;<br>
+ }<br>
+<br>
+ m_initialized = false;<br>
+}<br>
+<br>
+/* Try to load the cached compiled program binary, verify the device context is<br>
+ * still valid before reuse */<br>
+void ocl::cache_load(x265_param_t *param, char *devname, char *devvendor, char *driverversion)<br>
+{<br>
+ cl_int status;<br>
+<br>
+ /* try to load cached program binary */<br>
+ FILE *fp = fopen( param->pszCLbinFile, "rb" );<br>
+ if( !fp )<br>
+ return;<br>
+<br>
+ fseek( fp, 0L, SEEK_END );<br>
+ size_t size = ftell( fp );<br>
+ rewind( fp );<br>
+ uint8_t *binary = (uint8_t*)X265_MALLOC( uint8_t, size );<br>
+ if (!binary)<br>
+ return;<br>
+<br>
+ fread( binary, 1, size, fp );<br>
+ const uint8_t *ptr = (const uint8_t*)binary;<br>
+<br>
+#define CHECK_STRING( STR )\<br>
+ {\<br>
+ size_t len = strlen( STR );\<br>
+ if( size <= len || strncmp( (char*)ptr, STR, len ) )\<br>
+ goto fail;\<br>
+ else {\<br>
+ size -= (len+1); ptr += (len+1);\<br>
+ }\<br>
+ }<br>
+<br>
+ CHECK_STRING( devname );<br>
+ CHECK_STRING( devvendor );<br>
+ CHECK_STRING( driverversion );<br>
+ //CHECK_STRING( x265_opencl_source_hash );<br>
+#undef CHECK_STRING<br>
+<br>
+ m_program = clCreateProgramWithBinary( m_context, 1, &m_device, &size, &ptr, NULL, &status );<br>
+ if (status != CL_SUCCESS)<br>
+ {<br>
+ m_program = NULL;<br>
+ }<br>
+<br>
+fail:<br>
+ fclose( fp );<br>
+ X265_FREE( binary );<br>
+}<br>
+<br>
+/* Save the compiled program binary to a file for later reuse. Device context<br>
+ * is also saved in the cache file so we do not reuse stale binaries */<br>
+void ocl::cache_save(x265_param_t *param, char *devname, char *devvendor, char *driverversion)<br>
+{<br>
+ FILE *fp = fopen(param->pszCLbinFile, "wb");<br>
+ if( !fp )<br>
+ {<br>
+ x265_log(param, X265_LOG_INFO, "OpenCL: unable to open clbin file for write");<br>
+ return;<br>
+ }<br>
+<br>
+ size_t size;<br>
+ cl_int status = clGetProgramInfo( m_program, CL_PROGRAM_BINARY_SIZES, sizeof(size_t), &size, NULL );<br>
+ if( status == CL_SUCCESS )<br>
+ {<br>
+ uint8_t *binary = (uint8_t*)X265_MALLOC(uint8_t, size);<br>
+ status = clGetProgramInfo( m_program, CL_PROGRAM_BINARIES, sizeof(uint8_t *), &binary, NULL );<br>
+ if( status == CL_SUCCESS )<br>
+ {<br>
+ fputs( devname, fp );<br>
+ fputc( '\n', fp );<br>
+ fputs( devvendor, fp );<br>
+ fputc( '\n', fp );<br>
+ fputs( driverversion, fp );<br>
+ fputc( '\n', fp );<br>
+ //fputs( x265_opencl_source_hash, fp );<br>
+ //fputc( '\n', fp );<br>
+ fwrite( binary, 1, size, fp );<br>
+ }<br>
+ else<br>
+ x265_log(param, X265_LOG_INFO, "OpenCL: Unable to query program binary, no cache file generated");<br>
+ X265_FREE( binary );<br>
+ }<br>
+ else<br>
+ x265_log(param, X265_LOG_INFO, "OpenCL: Unable to query program binary size, no cache file generated");<br>
+ fclose( fp );<br>
+<br>
+ return;<br>
+}<br>
+<br>
+<br>
+oclME::oclME(x265_param_t *param)<br>
+ : m_pageLockedBuffer(NULL)<br>
+ , m_pageLockedPtr(NULL)<br>
+ , m_plOccupancy(0)<br>
+ , m_numCopies(0)<br>
+ , m_mvCostsBuffer(NULL)<br>
+ , m_mvpBuffer(NULL)<br>
+{<br>
+ for (int i = 0; i < 2; i++)<br>
+ {<br>
+ m_mvBuffer[i];<br>
+ m_costsBuffer[i];<br>
+ m_frameStatsBuffer[i];<br>
+ }<br>
+<br>
+ m_ocl = &g_ocl;<br>
+ init(param);<br>
+}<br>
+<br>
+/* The OpenCL source under common/opencl will be merged into common/oclkernels.h<br>
+ * by CMake. It defines a x265_opencl_source string which we will pass<br>
+ * to clCreateProgramWithSource(). We also attempt to use a cache file for the<br>
+ * compiled binary, stored in the current working folder. */<br>
+bool ocl::compile(x265_param_t *param)<br>
+{<br>
+ cl_int status;<br>
+ char devname[64];<br>
+ char devvendor[64];<br>
+ char driverversion[64];<br>
+<br>
+ status = clGetDeviceInfo( m_device, CL_DEVICE_NAME, sizeof(devname), devname, NULL );<br>
+ status |= clGetDeviceInfo( m_device, CL_DEVICE_VENDOR, sizeof(devvendor), devvendor, NULL );<br>
+ status |= clGetDeviceInfo( m_device, CL_DRIVER_VERSION, sizeof(driverversion), driverversion, NULL );<br>
+ if (status != CL_SUCCESS)<br>
+ {<br>
+ return false;<br>
+ }<br>
+<br>
+ // Most AMD GPUs have vector registers<br>
+ int vectorize = !strcmp( devvendor, "Advanced Micro Devices, Inc." );<br>
+ m_deviceAMDSI = false;<br>
+<br>
+ if (vectorize)<br>
+ {<br>
+ /* Disable OpenCL on Intel/AMD switchable graphics devices */<br>
+ if (detectSwitchableGraphics())<br>
+ {<br>
+ x265_log(param, X265_LOG_INFO, "OpenCL acceleration disabled, switchable graphics detected\n" );<br>
+ return false;<br>
+ }<br>
+<br>
+ /* Detect AMD SouthernIsland or newer device (single-width registers) */<br>
+ cl_uint simdwidth = 4;<br>
+ status = clGetDeviceInfo( m_device, CL_DEVICE_SIMD_INSTRUCTION_WIDTH_AMD, sizeof(cl_uint), &simdwidth, NULL );<br>
+ if( status == CL_SUCCESS && simdwidth == 1 )<br>
+ {<br>
+ vectorize = 0;<br>
+ m_deviceAMDSI = true;<br>
+ }<br>
+ }<br>
+<br>
+ x265_log(param, X265_LOG_INFO, "OpenCL acceleration enabled with %s %s %s\n", devvendor, devname, m_deviceAMDSI ? "(SI)" : "" );<br>
+<br>
+ cache_load(param, devname, devvendor, driverversion);<br>
+ if( !m_program )<br>
+ {<br>
+ /* clCreateProgramWithSource() requires a pointer variable, you cannot just use &x265_opencl_source */<br>
+ x265_log(param, X265_LOG_INFO, "Compiling OpenCL kernels...\n" );<br>
+ size_t size = strlen(x265_opencl_source);<br>
+ m_program = clCreateProgramWithSource( m_context, 1, &x265_opencl_source, &size, &status );<br>
+ if( status != CL_SUCCESS || !m_program )<br>
+ {<br>
+ x265_log(param, X265_LOG_WARNING, "OpenCL: unable to create program\n" );<br>
+ return false;<br>
+ }<br>
+ }<br>
+<br>
+ /* Build the program binary for the OpenCL device */<br>
+ const char *buildopts = vectorize ? "-DVECTORIZE=1" : "";<br>
+ status = clBuildProgram( m_program, 1, &m_device, buildopts, NULL, NULL );<br>
+ if( status == CL_SUCCESS )<br>
+ {<br>
+ cache_save(param, devname, devvendor, driverversion);<br>
+ return true;<br>
+ }<br>
+<br>
+ /* Compile failure, should not happen with production code. */<br>
+<br>
+ size_t build_log_len = 0;<br>
+<br>
+ status = clGetProgramBuildInfo( m_program, m_device, CL_PROGRAM_BUILD_LOG, build_log_len, NULL, &build_log_len );<br>
+ if( status != CL_SUCCESS )<br>
+ {<br>
+ x265_log(param, X265_LOG_WARNING, "OpenCL: Compilation failed, unable to query build log\n" );<br>
+ return false;<br>
+ }<br>
+<br>
+ char *build_log = (char*)X265_MALLOC(char, build_log_len);<br>
+ if( !build_log )<br>
+ {<br>
+ x265_log(param, X265_LOG_WARNING, "OpenCL: Compilation failed, unable to alloc build log\n" );<br>
+ return false;<br>
+ }<br>
+<br>
+ status = clGetProgramBuildInfo( m_program, m_device, CL_PROGRAM_BUILD_LOG, build_log_len, build_log, NULL );<br>
+ if( status != CL_SUCCESS )<br>
+ {<br>
+ x265_log(param, X265_LOG_WARNING, "OpenCL: Compilation failed, unable to get build log\n" );<br>
+ X265_FREE( build_log );<br>
+ return false;<br>
+ }<br>
+<br>
+ FILE *lg = fopen( "x265_kernel_build_log.txt", "w" );<br>
+ if( lg )<br>
+ {<br>
+ fwrite( build_log, 1, build_log_len, lg );<br>
+ fclose( lg );<br>
+ x265_log(param, X265_LOG_WARNING, "OpenCL: kernel build errors written to x265_kernel_build_log.txt\n" );<br>
+ }<br>
+<br>
+ X265_FREE( build_log );<br>
+<br>
+ clReleaseProgram(m_program);<br>
+ m_program = NULL;<br>
+<br>
+ return false;<br>
+}<br>
+<br>
+oclME::~oclME()<br>
+{<br>
+}<br>
+<br>
+void oclME::destroy()<br>
+{<br>
+}<br>
+<br>
+bool oclME::init(x265_param_t *param)<br>
+{<br>
+ cl_int status;<br>
+<br>
+ if (!m_ocl)<br>
+ {<br>
+ return false;<br>
+ }<br>
+ m_ocl->init(param);<br>
+<br>
+ m_pageLockedBuffer = clCreateBuffer(m_ocl->m_context, CL_MEM_WRITE_ONLY|CL_MEM_ALLOC_HOST_PTR, PAGE_LOCKED_BUF_SIZE, NULL, &status);<br>
+ if( status != CL_SUCCESS )<br>
+ {<br>
+ x265_log(param, X265_LOG_WARNING, "OpenCL: Unable to allocate page-locked buffer, error '%d'\n", status );<br>
+ return false;<br>
+ }<br>
+ m_pageLockedPtr = (char*) clEnqueueMapBuffer(m_ocl->m_queue, m_pageLockedBuffer, CL_TRUE, CL_MAP_READ | CL_MAP_WRITE,<br>
+ 0, PAGE_LOCKED_BUF_SIZE, 0, NULL, NULL, &status);<br>
+ if( status != CL_SUCCESS )<br>
+ {<br>
+ x265_log(param, X265_LOG_WARNING, "OpenCL: Unable to map page-locked buffer, error '%d'\n", status );<br>
+ return false;<br>
+ }<br>
+<br>
+ return true;<br>
+}<br></blockquote><div><br></div><div>let's not pull in all this ADL stuff until we have to</div><div> </div><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">
+/* OpenCL misbehaves on hybrid laptops with Intel iGPU and AMD dGPU, so<br>
+ * we consult AMD's ADL interface to detect this situation and disable<br>
+ * OpenCL on these machines (Linux and Windows) */<br>
+#ifndef _WIN32<br>
+#define __stdcall<br>
+#define HINSTANCE void *<br>
+#endif<br>
+typedef void* ( __stdcall *ADL_MAIN_MALLOC_CALLBACK )( int );<br>
+typedef int ( *ADL_MAIN_CONTROL_CREATE )(ADL_MAIN_MALLOC_CALLBACK, int );<br>
+typedef int ( *ADL_ADAPTER_NUMBEROFADAPTERS_GET ) ( int* );<br>
+typedef int ( *ADL_POWERXPRESS_SCHEME_GET ) ( int, int *, int *, int * );<br>
+typedef int ( *ADL_MAIN_CONTROL_DESTROY )();<br>
+#define ADL_OK 0<br>
+#define ADL_PX_SCHEME_DYNAMIC 2<br>
+<br>
+void* __stdcall adl_malloc_wrapper( int iSize ) { return X265_MALLOC(char, iSize); }<br>
+<br>
+int ocl::detectSwitchableGraphics()<br>
+{<br>
+ ADL_MAIN_CONTROL_CREATE ADL_Main_Control_Create;<br>
+ ADL_ADAPTER_NUMBEROFADAPTERS_GET ADL_Adapter_NumberOfAdapters_Get;<br>
+ ADL_POWERXPRESS_SCHEME_GET ADL_PowerXpress_Scheme_Get;<br>
+ ADL_MAIN_CONTROL_DESTROY ADL_Main_Control_Destroy;<br>
+ HINSTANCE hDLL;<br>
+ int ret = 0;<br>
+<br>
+#if _WIN32<br>
+ hDLL = LoadLibrary( "atiadlxx.dll" );<br>
+ if( !hDLL )<br>
+ hDLL = LoadLibrary( "atiadlxy.dll" );<br>
+#else<br>
+ hDLL = dlopen( "libatiadlxx.so", RTLD_LAZY|RTLD_GLOBAL );<br>
+#define GetProcAddress dlsym<br>
+#endif<br>
+ if( !hDLL )<br>
+ return ret;<br>
+<br>
+ ADL_Main_Control_Create = (ADL_MAIN_CONTROL_CREATE) GetProcAddress(hDLL, "ADL_Main_Control_Create");<br>
+ ADL_Main_Control_Destroy = (ADL_MAIN_CONTROL_DESTROY) GetProcAddress(hDLL, "ADL_Main_Control_Destroy");<br>
+ ADL_Adapter_NumberOfAdapters_Get = (ADL_ADAPTER_NUMBEROFADAPTERS_GET) GetProcAddress(hDLL, "ADL_Adapter_NumberOfAdapters_Get");<br>
+ ADL_PowerXpress_Scheme_Get = (ADL_POWERXPRESS_SCHEME_GET) GetProcAddress(hDLL, "ADL_PowerXpress_Scheme_Get");<br>
+ if( !ADL_Main_Control_Destroy || !ADL_Main_Control_Destroy || !ADL_Adapter_NumberOfAdapters_Get ||<br>
+ !ADL_PowerXpress_Scheme_Get )<br>
+ goto bail;<br>
+<br>
+ if( ADL_OK != ADL_Main_Control_Create( adl_malloc_wrapper, 1) )<br>
+ goto bail;<br>
+<br>
+ int numAdapters = 0;<br>
+ if( ADL_OK != ADL_Adapter_NumberOfAdapters_Get( &numAdapters ) )<br>
+ {<br>
+ ADL_Main_Control_Destroy();<br>
+ goto bail;<br>
+ }<br>
+<br>
+ for( int i = 0; i < numAdapters; i++ )<br>
+ {<br>
+ int PXSchemeRange, PXSchemeCurrentState, PXSchemeDefaultState;<br>
+ if( ADL_OK != ADL_PowerXpress_Scheme_Get( i, &PXSchemeRange, &PXSchemeCurrentState, &PXSchemeDefaultState) )<br>
+ break;<br>
+<br>
+ if( PXSchemeRange >= ADL_PX_SCHEME_DYNAMIC )<br>
+ {<br>
+ ret = 1;<br>
+ break;<br>
+ }<br>
+ }<br>
+<br>
+ ADL_Main_Control_Destroy();<br>
+<br>
+bail:<br>
+#if _WIN32<br>
+ FreeLibrary( hDLL );<br>
+#else<br>
+ dlclose( hDLL );<br>
+#endif<br>
+<br>
+ return ret;<br>
+}<br>
+<br>
diff -r 0d33ff236f68 -r 97e1bed019a6 source/common/opencl.h<br>
--- /dev/null Thu Jan 01 00:00:00 1970 +0000<br>
+++ b/source/common/opencl.h Tue Sep 17 13:41:27 2013 -0500<br>
@@ -0,0 +1,128 @@<br>
+/*****************************************************************************<br>
+ * opencl.h: OpenCL structures and defines<br>
+ *****************************************************************************<br>
+ * Copyright (C) 2013 x265 project<br>
+ *<br>
+ * Authors: Steve Borho <<a href="mailto:steve@borho.org">steve@borho.org</a>><br>
+ *<br>
+ * This program is free software; you can redistribute it and/or modify<br>
+ * it under the terms of the GNU General Public License as published by<br>
+ * the Free Software Foundation; either version 2 of the License, or<br>
+ * (at your option) any later version.<br>
+ *<br>
+ * This program is distributed in the hope that it will be useful,<br>
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of<br>
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the<br>
+ * GNU General Public License for more details.<br>
+ *<br>
+ * You should have received a copy of the GNU General Public License<br>
+ * along with this program; if not, write to the Free Software<br>
+ * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02111, USA.<br>
+ *<br>
+ * This program is also available under a commercial proprietary license.<br>
+ * For more information, contact us at <a href="mailto:licensing@multicorewareinc.com">licensing@multicorewareinc.com</a>.<br>
+ *****************************************************************************/<br>
+<br>
+#ifndef X265_OPENCL_H<br>
+#define X265_OPENCL_H<br>
+<br>
+#if defined __APPLE__<br>
+#include <OpenCL/opencl.h><br>
+#else<br>
+#include <CL/opencl.h><br>
+#endif<br>
+<br>
+#include "x265.h"<br>
+<br>
+namespace x265<br>
+{<br>
+<br>
+/* Number of PCIe copies that can be queued before requiring a flush */<br>
+#define MAX_FINISH_COPIES 1024<br>
+<br>
+/* Size (in bytes) of the page-locked buffer used for PCIe xfers */<br>
+#define PAGE_LOCKED_BUF_SIZE 32 * 1024 * 1024<br>
+<br>
+/* We define CL_QUEUE_THREAD_HANDLE_AMD here because it is not defined<br>
+ * in the OpenCL headers shipped with NVIDIA drivers. We need to be<br>
+ * able to compile on an NVIDIA machine and run optimally on an AMD GPU. */<br>
+#define CL_QUEUE_THREAD_HANDLE_AMD 0x403E<br>
+<br>
+#define OCLCHECK( method, ... )\<br>
+ status = method( __VA_ARGS__ );\<br>
+ if( status != CL_SUCCESS ) {\<br>
+ param->useOpenCL = 0;\<br>
+ m_ocl->m_fatalError = 1;\<br>
+ x264_log( param, X264_LOG_ERROR, # method " error '%d'\n", status );\<br>
+ return status;\<br>
+ }<br>
+<br>
+class ocl<br>
+{<br>
+public:<br>
+<br>
+ ocl();<br>
+ ~ocl();<br>
+<br>
+ cl_context m_context;<br>
+ cl_device_id m_device;<br>
+ cl_command_queue m_queue;<br>
+ cl_program m_program;<br>
+<br>
+ bool m_initialized;<br>
+ bool m_deviceAMDSI;<br>
+ bool m_fatalError;<br>
+ int m_openclThreadPri;<br>
+<br>
+ cl_kernel m_kernelMeStar;<br>
+<br>
+ bool init(x265_param_t *param);<br>
+ void destroy();<br>
+ int detectSwitchableGraphics();<br>
+<br>
+private:<br>
+<br>
+ bool createKernels(x265_param_t *param);<br>
+ bool compile(x265_param_t *param);<br>
+ void cache_load(x265_param_t *param, char *devname, char *devvendor, char *driverversion);<br>
+ void cache_save(x265_param_t *param, char *devname, char *devvendor, char *driverversion);<br>
+};<br>
+<br>
+<br>
+/* motion estimation */<br>
+class oclME<br>
+{<br>
+public:<br>
+<br>
+ oclME(x265_param_t *param);<br>
+ ~oclME();<br>
+<br>
+ ocl* m_ocl;<br>
+<br>
+ cl_mem m_pageLockedBuffer;<br>
+ char *m_pageLockedPtr;<br>
+ int m_plOccupancy;<br>
+<br>
+ struct<br>
+ {<br>
+ void *src;<br>
+ void *dest;<br>
+ int bytes;<br>
+ } m_copies[MAX_FINISH_COPIES];<br>
+ int m_numCopies;<br>
+<br>
+ cl_mem m_mvBuffer[2];<br>
+ cl_mem m_mvCostsBuffer;<br>
+ cl_mem m_mvpBuffer;<br>
+<br>
+ cl_mem m_costsBuffer[2];<br>
+ cl_mem m_frameStatsBuffer[2]; /* cost_est, cost_est_aq, intra_mbs */<br>
+<br>
+ bool init(x265_param_t *param);<br>
+ void destroy();<br>
+<br>
+};<br>
+<br>
+} // end namespace x265<br>
+<br>
+#endif<br>
diff -r 0d33ff236f68 -r 97e1bed019a6 source/common/opencl/<a href="http://me.cl" target="_blank">me.cl</a><br>
--- /dev/null Thu Jan 01 00:00:00 1970 +0000<br>
+++ b/source/common/opencl/<a href="http://me.cl" target="_blank">me.cl</a> Tue Sep 17 13:41:27 2013 -0500<br>
@@ -0,0 +1,15 @@<br>
+/* motion estimate */<br></blockquote><div><br></div><div>Oh interesting</div><div><br></div><div>It's pretty unlikely for a number of reasons that we'll be able to use the image (texture) data packing hack that we used in x264.</div>
<div> </div><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">
+__kernel void motion_estimate( read_only image2d_t fenc,<br>
+ read_only image2d_t fref,<br>
+ const __global short2 *in_mvs,<br>
+ __global short2 *out_mvs,<br>
+ __global short *out_mv_costs,<br>
+ __global short2 *mvp_buffer,<br>
+ __local short *cost_local,<br>
+ __local short2 *mvc_local,<br>
+ int mb_width,<br>
+ int lambda,<br>
+ int me_range)<br>
+{<br>
+}<br>
diff -r 0d33ff236f68 -r 97e1bed019a6 source/encoder/CMakeLists.txt<br>
--- a/source/encoder/CMakeLists.txt Mon Sep 16 21:06:49 2013 +0800<br>
+++ b/source/encoder/CMakeLists.txt Tue Sep 17 13:41:27 2013 -0500<br>
@@ -1,43 +1,47 @@<br></blockquote><div><br></div><div>another EOLN switch</div><div> </div><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">
-if(GCC)<br>
- # encoder.cpp must include HM headers which are not careful about named parameters<br>
- set_source_files_properties(encoder.cpp PROPERTIES COMPILE_FLAGS -Wno-unused-parameter)<br>
-endif(GCC)<br>
-<br>
-if(GCC)<br>
- if (NOT X64)<br>
- # force gcc to generate code for sync primitives<br>
- set_source_files_properties(framefilter.cpp PROPERTIES COMPILE_FLAGS -march=i686)<br>
- endif()<br>
-endif(GCC)<br>
-<br>
-file(GLOB LIBENCODER_HDR ../Lib/TLibEncoder/*.h)<br>
-file(GLOB LIBENCODER_SRC ../Lib/TLibEncoder/*.cpp)<br>
-source_group(TLibEncoder FILES ${LIBENCODER_SRC})<br>
-source_group(TLibEncoderH FILES ${LIBENCODER_HDR})<br>
-if(GCC)<br>
- set_source_files_properties(compress.cpp ${LIBENCODER_SRC} PROPERTIES COMPILE_FLAGS<br>
- "-Wno-sign-compare")<br>
-endif(GCC)<br>
-if(MSVC)<br>
- # ignore these warnings from HM source<br>
- # /wd4244 type conversion, possible loss of data<br>
- # /wd4512 assignment operator could not be generated<br>
- # /wd4127 conditional expression is constant<br>
- # /wd4389 signed/unsigned mismatch<br>
- # /wd4018 '<' signed/unsigned mismatch<br>
- # /wd4800 performance warning: bool coersion<br>
- set_source_files_properties(${LIBENCODER_SRC} PROPERTIES COMPILE_FLAGS<br>
- "/wd4244 /wd4512 /wd4127 /wd4389 /wd4018 /wd4800")<br>
-endif(MSVC)<br>
-<br>
-add_library(encoder STATIC ../../COPYING ../x265.h<br>
- ${LIBENCODER_SRC} ${LIBENCODER_HDR}<br>
- bitcost.cpp bitcost.h<br>
- motion.cpp motion.h<br>
- slicetype.cpp slicetype.h<br>
- frameencoder.cpp frameencoder.h<br>
- framefilter.cpp framefilter.h<br>
- cturow.cpp cturow.h<br>
- dpb.cpp dpb.h<br>
- ratecontrol.cpp ratecontrol.h<br>
- compress.cpp encoder.cpp)<br>
+if(GCC)<br>
+ # encoder.cpp must include HM headers which are not careful about named parameters<br>
+ set_source_files_properties(encoder.cpp PROPERTIES COMPILE_FLAGS -Wno-unused-parameter)<br>
+endif(GCC)<br>
+<br>
+if(GCC)<br>
+ if (NOT X64)<br>
+ # force gcc to generate code for sync primitives<br>
+ set_source_files_properties(framefilter.cpp PROPERTIES COMPILE_FLAGS -march=i686)<br>
+ endif()<br>
+endif(GCC)<br>
+<br>
+file(GLOB LIBENCODER_HDR ../Lib/TLibEncoder/*.h)<br>
+file(GLOB LIBENCODER_SRC ../Lib/TLibEncoder/*.cpp)<br>
+source_group(TLibEncoder FILES ${LIBENCODER_SRC})<br>
+source_group(TLibEncoderH FILES ${LIBENCODER_HDR})<br>
+if(GCC)<br>
+ set_source_files_properties(compress.cpp ${LIBENCODER_SRC} PROPERTIES COMPILE_FLAGS<br>
+ "-Wno-sign-compare")<br>
+endif(GCC)<br>
+if(MSVC)<br>
+ # ignore these warnings from HM source<br>
+ # /wd4244 type conversion, possible loss of data<br>
+ # /wd4512 assignment operator could not be generated<br>
+ # /wd4127 conditional expression is constant<br>
+ # /wd4389 signed/unsigned mismatch<br>
+ # /wd4018 '<' signed/unsigned mismatch<br>
+ # /wd4800 performance warning: bool coersion<br>
+ set_source_files_properties(${LIBENCODER_SRC} PROPERTIES COMPILE_FLAGS<br>
+ "/wd4244 /wd4512 /wd4127 /wd4389 /wd4018 /wd4800")<br>
+endif(MSVC)<br>
+<br>
+if(HAVE_OPENCL)<br>
+ include_directories(${OPENCL_INCLUDE_DIRS})<br>
+endif(HAVE_OPENCL)<br>
+<br>
+add_library(encoder STATIC ../../COPYING ../x265.h<br>
+ ${LIBENCODER_SRC} ${LIBENCODER_HDR}<br>
+ bitcost.cpp bitcost.h<br>
+ motion.cpp motion.h<br>
+ slicetype.cpp slicetype.h<br>
+ frameencoder.cpp frameencoder.h<br>
+ framefilter.cpp framefilter.h<br>
+ cturow.cpp cturow.h<br>
+ dpb.cpp dpb.h<br>
+ ratecontrol.cpp ratecontrol.h<br>
+ compress.cpp encoder.cpp)<br>
diff -r 0d33ff236f68 -r 97e1bed019a6 source/x265.h<br>
--- a/source/x265.h Mon Sep 16 21:06:49 2013 +0800<br>
+++ b/source/x265.h Tue Sep 17 13:41:27 2013 -0500<br>
@@ -277,6 +277,12 @@<br>
int qp; ///< Constant QP base value<br>
int rateFactor; ///< Constant rate factor (CRF)<br>
} rc;<br>
+<br>
+ // opencl<br>
+ int useOpenCL; ///< use OpenCL when available<br>
+ int iOpenCLDevice; ///< specify count of GPU devices to skip, for CLI users<br>
+ void *OpenCLDeviceId; ///< pass explicit cl_device_id as void*, for API users<br>
+ char *pszCLbinFile; ///< compiled OpenCL kernel cache file<br>
}<br>
x265_param_t;<br>
<br>
_______________________________________________<br>
x265-devel mailing list<br>
<a href="mailto:x265-devel@videolan.org">x265-devel@videolan.org</a><br>
<a href="https://mailman.videolan.org/listinfo/x265-devel" target="_blank">https://mailman.videolan.org/listinfo/x265-devel</a><br>
</blockquote></div><br><br clear="all"><div><br></div>-- <br>Steve Borho
</div></div>