[beignet] 04/07: Drop patches applied upstream
Rebecca Palmer
rnpalmer-guest at moszumanska.debian.org
Sun Aug 9 15:34:57 UTC 2015
This is an automated email from the git hooks/post-receive script.
rnpalmer-guest pushed a commit to branch master
in repository beignet.
commit 550051792ffc6b4416261d0221b9a59783fda6ea
Author: Rebecca N. Palmer <rebecca_palmer at zoho.com>
Date: Sun Aug 9 14:47:57 2015 +0100
Drop patches applied upstream
---
debian/changelog | 1 +
.../patches/builtin_pow-fix-spurious-failure.patch | 38 -------
debian/patches/default-to-full-precision.patch | 84 ---------------
debian/patches/disable-broken-fast-atomics.patch | 50 ---------
debian/patches/drop-structural_analysis.patch | 52 ---------
debian/patches/python3.patch | 101 -----------------
debian/patches/self-test.patch | 116 --------------------
debian/patches/series | 7 --
debian/patches/tgamma-accuracy.patch | 119 ---------------------
9 files changed, 1 insertion(+), 567 deletions(-)
diff --git a/debian/changelog b/debian/changelog
index ade126f..56eedda 100644
--- a/debian/changelog
+++ b/debian/changelog
@@ -2,6 +2,7 @@ beignet (1.1.0-1) UNRELEASED; urgency=medium
* New upstream release.
* Update and improve d/copyright.
+ * Drop patches applied upstream.
-- Rebecca N. Palmer <rebecca_palmer at zoho.com> Sun, 09 Aug 2015 13:27:07 +0100
diff --git a/debian/patches/builtin_pow-fix-spurious-failure.patch b/debian/patches/builtin_pow-fix-spurious-failure.patch
deleted file mode 100644
index cbec5da..0000000
--- a/debian/patches/builtin_pow-fix-spurious-failure.patch
+++ /dev/null
@@ -1,38 +0,0 @@
-Description: Declared lack of denormals is not a bug
-
-...so don't fail the test when the denormal 0.01**20.5 is flushed to 0
-
-Author: Rebecca N. Palmer <rebecca_palmer at zoho.com>
-Forwarded: http://lists.freedesktop.org/archives/beignet/2015-April/005549.html
-
---- beignet-1.0.2.orig/utests/builtin_pow.cpp
-+++ beignet-1.0.2/utests/builtin_pow.cpp
-@@ -38,6 +38,9 @@ static void builtin_pow(void)
- input_data2[i*count_input_ori+k] = ori_data[k];
- }
-
-+ cl_device_fp_config fp_config;
-+ clGetDeviceInfo(device, CL_DEVICE_SINGLE_FP_CONFIG, sizeof(cl_device_fp_config), &fp_config, 0);
-+ bool denormals_supported = fp_config & CL_FP_DENORM;
- const char* env_strict = getenv("OCL_STRICT_CONFORMANCE");
- float ULPSIZE_FACTOR = 16.0;
- if (env_strict != NULL && strcmp(env_strict, "0") == 0)
-@@ -76,7 +79,7 @@ static void builtin_pow(void)
- #if udebug
- if ( (isinf(cpu_data[index_cur]) && !isinf(gpu_data[index_cur])) ||
- (isnan(cpu_data[index_cur]) && !isnan(gpu_data[index_cur])) ||
-- (fabs(gpu_data[index_cur] - cpu_data[index_cur]) > cl_FLT_ULP(cpu_data[index_cur]) * ULPSIZE_FACTOR) )
-+ (fabs(gpu_data[index_cur] - cpu_data[index_cur]) > cl_FLT_ULP(cpu_data[index_cur]) * ULPSIZE_FACTOR && (denormals_supported || gpu_data[index_cur]!=0 || std::fpclassify(cpu_data[index_cur])!=FP_SUBNORMAL) ) )
- {
- printf_c("%d/%d: x:%f, y:%f -> gpu:%e cpu:%e err: %e ULP: %e\n", k, i, input_data1[k], input_data2[k], gpu_data[index_cur], cpu_data[index_cur],gpu_data[index_cur]-cpu_data[index_cur],cl_FLT_ULP(cpu_data[index_cur]));
- }
-@@ -89,7 +92,8 @@ static void builtin_pow(void)
- OCL_ASSERT(isnan(gpu_data[index_cur]));
- else
- {
-- OCL_ASSERT(fabs(gpu_data[index_cur] - cpu_data[index_cur]) < cl_FLT_ULP(cpu_data[index_cur]) * ULPSIZE_FACTOR);
-+ OCL_ASSERT((fabs(gpu_data[index_cur] - cpu_data[index_cur]) < cl_FLT_ULP(cpu_data[index_cur]) * ULPSIZE_FACTOR) ||
-+ (!denormals_supported && gpu_data[index_cur]==0 && std::fpclassify(cpu_data[index_cur])==FP_SUBNORMAL) );
- }
- #endif
- }
diff --git a/debian/patches/default-to-full-precision.patch b/debian/patches/default-to-full-precision.patch
deleted file mode 100644
index 96d2874..0000000
--- a/debian/patches/default-to-full-precision.patch
+++ /dev/null
@@ -1,84 +0,0 @@
-Description: Default to standard-compliant precision
-
-The OpenCL specification requires that any fast-but-imprecise mode
-be off by default: being too slow is usually more obvious than
-being too inaccurate.
-
-Author: Rebecca N. Palmer <rebecca_palmer at zoho.com>
-
---- beignet-1.0.2.orig/backend/src/backend/program.cpp
-+++ beignet-1.0.2/backend/src/backend/program.cpp
-@@ -113,7 +113,7 @@ namespace gbe {
-
- #ifdef GBE_COMPILER_AVAILABLE
- BVAR(OCL_OUTPUT_GEN_IR, false);
-- BVAR(OCL_STRICT_CONFORMANCE, false);
-+ BVAR(OCL_STRICT_CONFORMANCE, true);
-
- bool Program::buildFromLLVMFile(const char *fileName, const void* module, std::string &error, int optLevel) {
- ir::Unit *unit = new ir::Unit();
---- beignet-1.0.2.orig/docs/Beignet.mdwn
-+++ beignet-1.0.2/docs/Beignet.mdwn
-@@ -183,11 +183,10 @@ Known Issues
- * Precision issue.
- Currently Gen does not provide native support of high precision math functions
- required by OpenCL. We provide a software version to achieve high precision,
-- which you can turn on through
--
-- `# export OCL_STRICT_CONFORMANCE=1`.
--
-- But be careful, this would make your CL kernel run a little longer.
-+ controlled by the OCL_STRICT_CONFORMANCE environment variable. Default is
-+ 1 (high-precision) in Debian beignet but 0 (high-speed) in upstream beignet.
-+
-+ The native_* functions are always high-speed/low-precision.
-
- * cl\_khr\_gl\_sharing.
- This extension highly depends on mesa support. It seems that mesa would not provide
---- beignet-1.0.2.orig/docs/Beignet/Backend.mdwn
-+++ beignet-1.0.2/docs/Beignet/Backend.mdwn
-@@ -37,9 +37,8 @@ Environment variables are used all over
- precision math instructions compliant with OpenCL Spec. So we provide a
- software version to meet the high precision requirement. Obviously the
- software version's performance is not as good as native version supported by
-- GEN hardware. What's more, most graphics application don't need this high
-- precision, so we choose 0 as the default value. So OpenCL apps do not suffer
-- the performance penalty for using high precision math functions.
-+ GEN hardware. Default is 1 (high-precision) in Debian beignet but
-+ 0 (high-speed) in upstream beignet.
-
- - `OCL_SIMD_WIDTH` `(8 or 16)`. Select the number of lanes per hardware thread,
- Normally, you don't need to set it, we will select suitable simd width for
---- beignet-1.0.2.orig/utests/builtin_pow.cpp
-+++ beignet-1.0.2/utests/builtin_pow.cpp
-@@ -40,7 +40,7 @@ static void builtin_pow(void)
-
- const char* env_strict = getenv("OCL_STRICT_CONFORMANCE");
- float ULPSIZE_FACTOR = 16.0;
-- if (env_strict == NULL || strcmp(env_strict, "0") == 0)
-+ if (env_strict != NULL && strcmp(env_strict, "0") == 0)
- ULPSIZE_FACTOR = 10000.;
-
- OCL_CREATE_KERNEL("builtin_pow");
---- beignet-1.0.2.orig/utests/builtin_tgamma.cpp
-+++ beignet-1.0.2/utests/builtin_tgamma.cpp
-@@ -17,7 +17,7 @@ void builtin_tgamma(void)
- locals[0] = 16;
- const char* env_strict = getenv("OCL_STRICT_CONFORMANCE");
- float ULPSIZE_FACTOR = 16.0;
-- if (env_strict == NULL || strcmp(env_strict, "0") == 0)
-+ if (env_strict != NULL && strcmp(env_strict, "0") == 0)
- ULPSIZE_FACTOR = 10000.;
-
- for (int j = 0; j < 1024; j ++) {
---- beignet-1.0.2.orig/utests/utest_generator.py
-+++ beignet-1.0.2/utests/utest_generator.py
-@@ -109,7 +109,7 @@ def udebug(ulpSize,returnType,function):
-
- const char* env_strict = getenv("OCL_STRICT_CONFORMANCE");
-
-- if (env_strict == NULL || strcmp(env_strict, "0") == 0)
-+ if (env_strict != NULL && strcmp(env_strict, "0") == 0)
- ULPSIZE_FACTOR = 1000;
- else
- ULPSIZE_FACTOR = %s;
diff --git a/debian/patches/disable-broken-fast-atomics.patch b/debian/patches/disable-broken-fast-atomics.patch
deleted file mode 100644
index 61f0b9e..0000000
--- a/debian/patches/disable-broken-fast-atomics.patch
+++ /dev/null
@@ -1,50 +0,0 @@
-Description: Disable fast atomics, they break Haswell
-
-Revert "CL/Driver: enable atomics in L3 for HSW."
-(upstream ef7127c03bd533277afc443b335c37a69927250a)
-
-This optimization is rejected by a Linux security mechanism,
-making all kernels fail to run on Haswell with
-"drm_intel_gem_bo_context_exec() failed: Invalid argument"
-
-Origin: upstream 83f8739b6fc4893fac60145326052ccb5cf653dc
-Author: Zhigang Gong <zhigang.gong at intel.com>
-
-diff --git a/src/intel/intel_defines.h b/src/intel/intel_defines.h
-index 651e285..1080a91 100644
---- a/src/intel/intel_defines.h
-+++ b/src/intel/intel_defines.h
-@@ -305,10 +305,6 @@
- #define URB_SIZE(intel) (IS_IGDNG(intel->device_id) ? 1024 : \
- IS_G4X(intel->device_id) ? 384 : 256)
-
--// HSW
--#define HSW_SCRATCH1_OFFSET (0xB038)
--#define HSW_ROW_CHICKEN3_HDC_OFFSET (0xE49C)
--
- // L3 cache stuff
- #define GEN7_L3_SQC_REG1_ADDRESS_OFFSET (0XB010)
- #define GEN7_L3_CNTL_REG2_ADDRESS_OFFSET (0xB020)
-diff --git a/src/intel/intel_gpgpu.c b/src/intel/intel_gpgpu.c
-index ab4545e..b083dab 100644
---- a/src/intel/intel_gpgpu.c
-+++ b/src/intel/intel_gpgpu.c
-@@ -719,16 +719,7 @@ static void
- intel_gpgpu_set_L3_gen75(intel_gpgpu_t *gpgpu, uint32_t use_slm)
- {
- /* still set L3 in batch buffer for fulsim. */
-- BEGIN_BATCH(gpgpu->batch, 15);
-- OUT_BATCH(gpgpu->batch, CMD_LOAD_REGISTER_IMM | 1); /* length - 2 */
-- /* FIXME: KMD always disable the atomic in L3 for some reason.
-- I checked the spec, and don't think we need that workaround now.
-- Before I send a patch to kernel, let's just enable it here. */
-- OUT_BATCH(gpgpu->batch, HSW_SCRATCH1_OFFSET);
-- OUT_BATCH(gpgpu->batch, 0); /* enable atomic in L3 */
-- OUT_BATCH(gpgpu->batch, CMD_LOAD_REGISTER_IMM | 1); /* length - 2 */
-- OUT_BATCH(gpgpu->batch, HSW_ROW_CHICKEN3_HDC_OFFSET);
-- OUT_BATCH(gpgpu->batch, (1 << 6ul) << 16); /* enable atomic in L3 */
-+ BEGIN_BATCH(gpgpu->batch, 9);
- OUT_BATCH(gpgpu->batch, CMD_LOAD_REGISTER_IMM | 1); /* length - 2 */
- OUT_BATCH(gpgpu->batch, GEN7_L3_SQC_REG1_ADDRESS_OFFSET);
- OUT_BATCH(gpgpu->batch, 0x08800000);
-
diff --git a/debian/patches/drop-structural_analysis.patch b/debian/patches/drop-structural_analysis.patch
deleted file mode 100644
index af8e0be..0000000
--- a/debian/patches/drop-structural_analysis.patch
+++ /dev/null
@@ -1,52 +0,0 @@
-Description: Stop using non-DFSG gpuocelot code
-
-backend/src/ir/structural_analysis.* (derived from gpuocelot) have
-been deleted because they are not DFSG-free:
-https://lists.debian.org/debian-legal/2015/05/msg00008.html
-
-Origin: upstream 347116ef2855e73c25dc09c0a8ee1c7f58c46099
-Author: Zhigang Gong <zhigang.gong at intel.com>
-
-diff --git a/backend/src/CMakeLists.txt b/backend/src/CMakeLists.txt
-index a6736ec..45f18d8 100644
---- a/backend/src/CMakeLists.txt
-+++ b/backend/src/CMakeLists.txt
-@@ -66,8 +66,6 @@ set (GBE_SRC
- ir/lowering.hpp
- ir/printf.cpp
- ir/printf.hpp
-- ir/structural_analysis.cpp
-- ir/structural_analysis.hpp
- ir/immediate.hpp
- ir/immediate.cpp
- backend/context.cpp
-diff --git a/backend/src/llvm/llvm_to_gen.cpp b/backend/src/llvm/llvm_to_gen.cpp
-index a4ce4a2..3e11fdb 100644
---- a/backend/src/llvm/llvm_to_gen.cpp
-+++ b/backend/src/llvm/llvm_to_gen.cpp
-@@ -61,7 +61,6 @@
- #include "sys/cvar.hpp"
- #include "sys/platform.hpp"
- #include "ir/unit.hpp"
--#include "ir/structural_analysis.hpp"
-
- #include <clang/CodeGen/CodeGenAction.h>
-
-@@ -308,16 +307,6 @@ namespace gbe
- // Print the code extra optimization passes
- OUTPUT_BITCODE(AFTER_GEN, mod);
-
-- const ir::Unit::FunctionSet& fs = unit.getFunctionSet();
-- ir::Unit::FunctionSet::const_iterator iter = fs.begin();
-- while(iter != fs.end())
-- {
-- analysis::ControlTree *ct = new analysis::ControlTree(iter->second);
-- ct->analyze();
-- delete ct;
-- iter++;
-- }
--
- delete libraryInfo;
- return true;
- }
-
diff --git a/debian/patches/python3.patch b/debian/patches/python3.patch
deleted file mode 100644
index f280114..0000000
--- a/debian/patches/python3.patch
+++ /dev/null
@@ -1,101 +0,0 @@
-Description: Allow building with Python 3
-
-Author: Rebecca N. Palmer <rebecca_palmer at zoho.com>
-Forwarded: http://lists.freedesktop.org/archives/beignet/2015-April/005584.html
-
---- beignet-1.0.3.orig/backend/src/libocl/script/gen_vector.py
-+++ beignet-1.0.3/backend/src/libocl/script/gen_vector.py
-@@ -20,13 +20,14 @@
-
- # This file is to generate inline code to lower down those builtin
- # vector functions to scalar functions.
-+from __future__ import print_function
- import re
- import sys
- import os
-
- if len(sys.argv) != 4:
-- print "Invalid argument {0}".format(sys.argv)
-- print "use {0} spec_file_name output_file_name just_proto".format(sys.argv[0])
-+ print("Invalid argument {0}".format(sys.argv))
-+ print("use {0} spec_file_name output_file_name just_proto".format(sys.argv[0]))
- raise
-
- all_vector = 1,2,3,4,8,16
-@@ -61,8 +62,8 @@ all_type = all_int_type + all_float_type
-
- # all vector/scalar types
- for t in all_type:
-- exec "{0}n = [\"{0}n\", gen_vector_type([\"{0}\"])]".format(t)
-- exec "s{0} = [\"{0}\", gen_vector_type([\"{0}\"], [1])]".format(t)
-+ exec("{0}n = [\"{0}n\", gen_vector_type([\"{0}\"])]".format(t))
-+ exec("s{0} = [\"{0}\", gen_vector_type([\"{0}\"], [1])]".format(t))
-
- # Predefined type sets according to the Open CL spec.
- math_gentype = ["math_gentype", gen_vector_type(all_float_type)]
-@@ -124,8 +125,8 @@ def check_type(types):
- for t in types:
- memspace, t = stripMemSpace(t)
- if not t in type_dict:
-- print t
-- raise "found invalid type."
-+ print(t)
-+ raise TypeError("found invalid type.")
-
- def match_unsigned(dtype):
- if dtype[0] == 'float':
-@@ -187,8 +188,8 @@ def fixup_type(dstType, srcType, n):
- if (len(dstType) == len(srcType)):
- return dstType[n]
-
-- print dstType, srcType
-- raise "type mispatch"
-+ print(dstType, srcType)
-+ raise TypeError("type mispatch")
-
- class builtinProto():
- valueTypeStr = ""
-@@ -226,7 +227,7 @@ class builtinProto():
-
- def init_from_line(self, t):
- self.append('//{0}'.format(t))
-- line = filter(None, re.split(',| |\(', t.rstrip(')\n')))
-+ line = [_f for _f in re.split(',| |\(', t.rstrip(')\n')) if _f]
- self.paramCount = 0
- stripped = 0
- memSpace = ''
-@@ -310,7 +311,7 @@ class builtinProto():
- vtype = fixup_type(vtypeSeq, ptypeSeqs[n], i)
- if vtype[1] != ptype[1]:
- if ptype[1] != 1:
-- raise "parameter is not a scalar but has different width with result value."
-+ raise TypeError("parameter is not a scalar but has different width with result value.")
- if isPointer(ptype):
- formatStr += '&'
- formatStr += 'param{0}'.format(n)
-@@ -333,7 +334,7 @@ class builtinProto():
-
- def output(self):
- for line in self.outputStr:
-- print line
-+ print(line)
-
- def output(self, outFile):
- for line in self.outputStr:
---- beignet-1.0.3.orig/utests/utest_generator.py
-+++ beignet-1.0.3/utests/utest_generator.py
-@@ -1,4 +1,5 @@
- #!/usr/bin/python
-+from __future__ import print_function
- import os,sys,re
-
- FLT_MAX_POSI='0x1.fffffep127f'
-@@ -327,7 +328,7 @@ which can print more values and informat
- file_object.close()
-
- def nameForCmake(self,content,namesuffix):
-- print("generated/%s_%s.cpp"%(self.fileName,namesuffix)),
-+ print("generated/%s_%s.cpp"%(self.fileName,namesuffix),end=" ")
-
- def utestFunc(self,index):
- funcLines=[]
diff --git a/debian/patches/self-test.patch b/debian/patches/self-test.patch
deleted file mode 100644
index b6ce93f..0000000
--- a/debian/patches/self-test.patch
+++ /dev/null
@@ -1,116 +0,0 @@
-Description: Test device in clGetDeviceIDs
-
-Run a small kernel to check that the device works (including the
-__local memory space, known to be problematic on Haswell), and
-hide non-working devices
-
-Author: Rebecca N. Palmer <rebecca_palmer at zoho.com>
-Forwarded: http://cgit.freedesktop.org/beignet/commit/?id=dfcc554d3bbc3bcc0810b3da5ee26b4345e6b4f7
-
---- beignet-1.0.3.orig/src/cl_device_id.c
-+++ beignet-1.0.3/src/cl_device_id.c
-@@ -458,6 +458,72 @@ brw_gt3_break:
- return ret;
- }
-
-+/* Runs a small kernel to check that the device works; returns
-+ * 0 for success, 1 for silently wrong result, 2 for error */
-+LOCAL cl_int
-+cl_self_test(cl_device_id device)
-+{
-+ cl_int status, ret;
-+ cl_context ctx;
-+ cl_command_queue queue;
-+ cl_program program;
-+ cl_kernel kernel;
-+ cl_mem buffer;
-+ cl_event kernel_finished;
-+ size_t n = 3;
-+ cl_int test_data[3] = {3, 7, 5};
-+ const char* kernel_source = "__kernel void self_test(__global int *buf) {"
-+ " __local int tmp[3];"
-+ " tmp[get_local_id(0)] = buf[get_local_id(0)];"
-+ " barrier(CLK_LOCAL_MEM_FENCE);"
-+ " buf[get_global_id(0)] = tmp[2 - get_local_id(0)] + buf[get_global_id(0)];"
-+ "}"; // using __local to catch the "no SLM on Haswell" problem
-+ ret = 2;
-+ ctx = clCreateContext(NULL, 1, &device, NULL, NULL, &status);
-+ if (status == CL_SUCCESS) {
-+ queue = clCreateCommandQueue(ctx, device, 0, &status);
-+ if (status == CL_SUCCESS) {
-+ program = clCreateProgramWithSource(ctx, 1, &kernel_source, NULL, &status);
-+ if (status == CL_SUCCESS) {
-+ status = clBuildProgram(program, 1, &device, "", NULL, NULL);
-+ if (status == CL_SUCCESS) {
-+ kernel = clCreateKernel(program, "self_test", &status);
-+ if (status == CL_SUCCESS) {
-+ buffer = clCreateBuffer(ctx, CL_MEM_COPY_HOST_PTR, n*4, test_data, &status);
-+ if (status == CL_SUCCESS) {
-+ status = clSetKernelArg(kernel, 0, sizeof(cl_mem), &buffer);
-+ if (status == CL_SUCCESS) {
-+ status = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &n, &n, 0, NULL, &kernel_finished);
-+ if (status == CL_SUCCESS) {
-+ status = clEnqueueReadBuffer(queue, buffer, CL_TRUE, 0, n*4, test_data, 1, &kernel_finished, NULL);
-+ if (status == CL_SUCCESS) {
-+ if (test_data[0] == 8 && test_data[1] == 14 && test_data[2] == 8){
-+ ret = 0;
-+ } else {
-+ ret = 1;
-+ printf("Beignet: self-test failed: (3, 7, 5) + (5, 7, 3) returned (%i, %i, %i)\n"
-+ "This is a known bug on Haswell systems, see /usr/share/doc/beignet-opencl-icd/README.Debian\n", test_data[0], test_data[1], test_data[2]);
-+ }
-+ }
-+ }
-+ }
-+ }
-+ clReleaseMemObject(buffer);
-+ }
-+ clReleaseKernel(kernel);
-+ }
-+ }
-+ clReleaseProgram(program);
-+ }
-+ clReleaseCommandQueue(queue);
-+ }
-+ clReleaseContext(ctx);
-+ if (ret == 2) {
-+ printf("Beignet: self-test failed: error %i\nSee /usr/share/doc/beignet-opencl-icd/README.Debian\n", status);
-+ }
-+ return ret;
-+}
-+
- LOCAL cl_int
- cl_get_device_ids(cl_platform_id platform,
- cl_device_type device_type,
-@@ -469,6 +535,20 @@ cl_get_device_ids(cl_platform_id plat
-
- /* Do we have a usable device? */
- device = cl_get_gt_device();
-+ if (device && cl_self_test(device)) {
-+ int disable_self_test = 0;
-+ // can't use BVAR (backend/src/sys/cvar.hpp) here as it's C++
-+ const char *env = getenv("OCL_IGNORE_SELF_TEST");
-+ if (env != NULL) {
-+ sscanf(env, "%i", &disable_self_test);
-+ }
-+ if (disable_self_test) {
-+ printf("Beignet: Warning - overriding self-test failure\n");
-+ } else {
-+ printf("Beignet: disabling non-working device\n");
-+ device = 0;
-+ }
-+ }
- if (!device) {
- if (num_devices)
- *num_devices = 0;
---- beignet-1.0.3.orig/utests/setenv.sh.in
-+++ beignet-1.0.3/utests/setenv.sh.in
-@@ -6,6 +6,8 @@ export OCL_PCH_PATH=@LOCAL_OCL_PCH_OBJEC
- export OCL_KERNEL_PATH=@CMAKE_CURRENT_SOURCE_DIR@/../kernels
- export OCL_GBE_PATH=@LOCAL_GBE_OBJECT_DIR@
- export OCL_INTERP_PATH=@LOCAL_INTERP_OBJECT_DIR@
-+#disable self-test so we can get something more precise than "doesn't work"
-+export OCL_IGNORE_SELF_TEST=1
- obj-`dpkg-architecture -qDEB_HOST_MULTIARCH`/utests/utest_run -a
- export OCL_STRICT_CONFORMANCE=0
- obj-`dpkg-architecture -qDEB_HOST_MULTIARCH`/utests/utest_run -a
diff --git a/debian/patches/series b/debian/patches/series
index 95f292b..a37a907 100644
--- a/debian/patches/series
+++ b/debian/patches/series
@@ -2,13 +2,6 @@ Debian-compliant-compiler-flags-handling.patch
Enable-test-debug.patch
support-kfreebsd.patch
reduce-notfound-output.patch
-default-to-full-precision.patch
shared-llvm.patch
-builtin_pow-fix-spurious-failure.patch
-tgamma-accuracy.patch
-python3.patch
-self-test.patch
-disable-broken-fast-atomics.patch
update-docs.patch
-drop-structural_analysis.patch
allow-+-in-directory.patch
diff --git a/debian/patches/tgamma-accuracy.patch b/debian/patches/tgamma-accuracy.patch
deleted file mode 100644
index f9b9b03..0000000
--- a/debian/patches/tgamma-accuracy.patch
+++ /dev/null
@@ -1,119 +0,0 @@
-Description: Make tgamma meet the OpenCL accuracy standard
-
-tgamma=exp(lgamma) was not accurate enough for approx. x>8
-
-Author: Rebecca N. Palmer <rebecca_palmer at zoho.com>
-Forwarded: http://lists.freedesktop.org/archives/beignet/2015-April/005548.html
-
---- beignet-1.0.2.orig/backend/src/libocl/tmpl/ocl_math.tmpl.cl
-+++ beignet-1.0.2/backend/src/libocl/tmpl/ocl_math.tmpl.cl
-@@ -1746,12 +1746,6 @@ OVERLOADABLE float __gen_ocl_internal_ex
- }
- }
-
--INLINE_OVERLOADABLE float tgamma(float x) {
-- float y;
-- int s;
-- y=lgamma_r(x,&s);
-- return __gen_ocl_internal_exp(y)*s;
--}
-
- /* erf,erfc from glibc s_erff.c -- float version of s_erf.c.
- * Conversion to float by Ian Lance Taylor, Cygnus Support, ian at cygnus.com.
-@@ -3167,6 +3167,95 @@ OVERLOADABLE float __gen_ocl_internal_po
- return sn*z;
- }
-
-+OVERLOADABLE float tgamma (float x)
-+{/*based on glibc tgammaf*/
-+ unsigned int hx;
-+
-+ GEN_OCL_GET_FLOAT_WORD(hx,x);
-+
-+ if (hx == 0xff800000)
-+ {
-+ /* x == -Inf. According to ISO this is NaN. */
-+ return NAN;
-+ }
-+ if ((hx & 0x7f800000) == 0x7f800000)
-+ {
-+ /* Positive infinity (return positive infinity) or NaN (return
-+ NaN). */
-+ return x;
-+ }
-+ if (x < 0.0f && __gen_ocl_internal_floor (x) == x)
-+ {
-+ /* integer x < 0 */
-+ return NAN;
-+ }
-+
-+ if (x >= 36.0f)
-+ {
-+ /* Overflow. */
-+ return INFINITY;
-+ }
-+ else if (x <= 0.0f && x >= -FLT_EPSILON / 4.0f)
-+ {
-+ return 1.0f / x;
-+ }
-+ else
-+ {
-+ float sinpix = __gen_ocl_internal_sinpi(x);
-+ if (x <= -42.0f)
-+ /* Underflow. */
-+ {return 0.0f * sinpix /*for sign*/;}
-+ int exp2_adj = 0;
-+ float x_abs = __gen_ocl_fabs(x);
-+ float gam0;
-+
-+ if (x_abs < 4.0f) {
-+ /* gamma = exp(lgamma) is only accurate for small lgamma */
-+ float prod,x_adj;
-+ if (x_abs < 0.5f) {
-+ prod = 1.0f / x_abs;
-+ x_adj = x_abs + 1.0f;
-+ } else if (x_abs <= 1.5f) {
-+ prod = 1.0f;
-+ x_adj = x_abs;
-+ } else if (x_abs < 2.5f) {
-+ x_adj = x_abs - 1.0f;
-+ prod = x_adj;
-+ } else {
-+ x_adj = x_abs - 2.0f;
-+ prod = x_adj * (x_abs - 1.0f);
-+ }
-+ gam0 = __gen_ocl_internal_exp (lgamma (x_adj)) * prod;
-+ }
-+ else {
-+ /* Compute gamma (X) using Stirling's approximation,
-+ starting by computing pow (X, X) with a power of 2
-+ factored out to avoid intermediate overflow. */
-+ float x_int = __gen_ocl_internal_round (x_abs);
-+ float x_frac = x_abs - x_int;
-+ int x_log2;
-+ float x_mant = frexp (x_abs, &x_log2);
-+ if (x_mant < M_SQRT1_2_F)
-+ {
-+ x_log2--;
-+ x_mant *= 2.0f;
-+ }
-+ exp2_adj = x_log2 * (int) x_int;
-+ float ret = (__gen_ocl_internal_pow(x_mant, x_abs)
-+ * exp2 (x_log2 * x_frac)
-+ * __gen_ocl_internal_exp (-x_abs)
-+ * sqrt (2.0f * M_PI_F / x_abs) );
-+
-+ float x2 = x_abs * x_abs;
-+ float bsum = (0x3.403404p-12f / x2 -0xb.60b61p-12f) / x2 + 0x1.555556p-4f;
-+ gam0 = ret + ret * __gen_ocl_internal_expm1 (bsum / x_abs);
-+ }
-+ if (x > 0.0f) {return __gen_ocl_internal_ldexp (gam0, exp2_adj);}
-+ float gam1 = M_PI_F / (-x * sinpix * gam0);
-+ return __gen_ocl_internal_ldexp (gam1, -exp2_adj);
-+ }
-+}
-+
- OVERLOADABLE float hypot(float x, float y) {
- if (__ocl_math_fastpath_flag)
- return __gen_ocl_internal_fastpath_hypot(x, y);
-
--
Alioth's /usr/local/bin/git-commit-notice on /srv/git.debian.org/git/pkg-opencl/beignet.git
More information about the Pkg-opencl-commits
mailing list