diff --git a/CMakeLists.txt b/CMakeLists.txt index cd6d9bdd..5f43b291 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -193,6 +193,8 @@ add_library(oclgrind ${CORE_LIB_TYPE} src/core/WorkItem.cpp src/core/WorkItemBuiltins.cpp src/core/WorkGroup.cpp + src/plugins/ArithmeticExceptions.h + src/plugins/ArithmeticExceptions.cpp src/plugins/InstructionCounter.h src/plugins/InstructionCounter.cpp src/plugins/InteractiveDebugger.h diff --git a/Makefile.am b/Makefile.am index 9ac71985..18681a33 100644 --- a/Makefile.am +++ b/Makefile.am @@ -43,6 +43,7 @@ liboclgrind_la_SOURCES = src/core/common.h src/core/common.cpp \ src/core/Queue.h src/core/Queue.cpp src/core/WorkItem.h \ src/core/WorkItem.cpp src/core/WorkItemBuiltins.cpp \ src/core/WorkGroup.h src/core/WorkGroup.cpp \ + src/plugins/ArithmeticExceptions.h src/plugins/ArithmeticExceptions.cpp \ src/plugins/InstructionCounter.h src/plugins/InstructionCounter.cpp \ src/plugins/InteractiveDebugger.h src/plugins/InteractiveDebugger.cpp \ src/plugins/Logger.h src/plugins/Logger.cpp src/plugins/MemCheck.h \ diff --git a/src/core/Context.cpp b/src/core/Context.cpp index e6fc4151..b7c5ea5a 100644 --- a/src/core/Context.cpp +++ b/src/core/Context.cpp @@ -36,6 +36,7 @@ #include "plugins/MemCheck.h" #include "plugins/RaceDetector.h" #include "plugins/Uninitialized.h" +#include "plugins/ArithmeticExceptions.h" using namespace oclgrind; using namespace std; @@ -97,6 +98,9 @@ void Context::loadPlugins() if (checkEnv("OCLGRIND_INTERACTIVE")) m_plugins.push_back(make_pair(new InteractiveDebugger(this), true)); + if (checkEnv("OCLGRIND_ARITHMETIC_EXCEPTIONS")) + m_plugins.push_back(make_pair(new ArithmeticExceptions(this), true)); + // Load dynamic plugins const char *dynamicPlugins = getenv("OCLGRIND_PLUGINS"); diff --git a/src/kernel/oclgrind-kernel.cpp b/src/kernel/oclgrind-kernel.cpp index 1ce599d6..93fc26ab 100644 --- a/src/kernel/oclgrind-kernel.cpp +++ b/src/kernel/oclgrind-kernel.cpp @@ -48,7 +48,11 @@ static bool parseArguments(int argc, char *argv[]) { for (int i = 1; i < argc; i++) { - if (!strcmp(argv[i], "--build-options")) + if (!strcmp(argv[i], "--arithmetic-exceptions")) + { + setEnvironment("OCLGRIND_ARITHMETIC_EXCEPTIONS", "1"); + } + else if (!strcmp(argv[i], "--build-options")) { if (++i >= argc) { @@ -190,6 +194,8 @@ static void printUsage() << " oclgrind-kernel [--help | --version]" << endl << endl << "Options:" << endl + << " --arithmetic-exceptions " + "Enable detection of arithmetic exceptions" << endl << " --build-options OPTIONS " "Additional options to pass to the OpenCL compiler" << endl << " --data-races " diff --git a/src/plugins/ArithmeticExceptions.cpp b/src/plugins/ArithmeticExceptions.cpp new file mode 100644 index 00000000..fa5bb571 --- /dev/null +++ b/src/plugins/ArithmeticExceptions.cpp @@ -0,0 +1,366 @@ +// ArithmeticExceptions.h (Oclgrind) +// Copyright (c) 2016, Moritz Pflanzer, +// Imperial College London. All rights reserved. +// +// This program is provided under a three-clause BSD license. For full +// license terms please see the LICENSE file distributed with this +// source code. + +// The recommendations how to detect undefined behaviour haven been adopted from: +// https://www.securecoding.cert.org/confluence/display/c/INT32-C.+Ensure+that+operations+on+signed+integers+do+not+result+in+overflow +// https://www.securecoding.cert.org/confluence/display/c/INT33-C.+Ensure+that+division+and+remainder+operations+do+not+result+in+divide-by-zero+errors +// Overflow in unary minus operator cannot be checked in LLVM IR +// Oversize shifts are well-defined in OpenCL +// Arithmetic operations involving vector types do not seem to produce nsw or nuw flags. This means overflows cannot be detected. + +#include "core/common.h" +#include "core/Context.h" +#include "core/WorkItem.h" + +#include +#include + +#include "llvm/IR/Instruction.h" +#include "llvm/IR/Instructions.h" +#include "llvm/IR/Type.h" + +#include "ArithmeticExceptions.h" + +#include "core/Kernel.h" +#include "core/KernelInvocation.h" + +using namespace oclgrind; +using namespace std; + +#define HALF_MAX 65504 + +static long getSignedMinValue(const unsigned int bits) +{ + return -(1L << (bits - 2)) - (1L << (bits - 2)); +} + +static long getSignedMaxValue(const unsigned int bits) +{ + return ((1L << (bits - 2)) - 1L) + (1L << (bits - 2)); +} + +static long getUnsignedMaxValue(const unsigned int bits) +{ + return ((1L << (bits - 1)) - 1L) + (1L << (bits - 1)); +} + +void ArithmeticExceptions::instructionExecuted( + const WorkItem *workItem, const llvm::Instruction *instruction, + const TypedValue& result) +{ + switch(instruction->getOpcode()) + { + case llvm::Instruction::Add: + { + const auto&& BinOp = llvm::dyn_cast(instruction); + + // Check for signed overflow + // Report an exception iff LLVM would create a poisoned value on overflow + // Other than that it is not possible to differentiate between signed/unsigned values + if(BinOp->hasNoSignedWrap()) + { + const TypedValue LHS = workItem->getOperand(instruction->getOperand(0)); + const TypedValue RHS = workItem->getOperand(instruction->getOperand(1)); + const unsigned int ResultWidth = BinOp->getType()->getScalarSizeInBits(); + + for(unsigned int i = 0; i < result.num; ++i) + { + if(((RHS.getSInt(i) > 0) && (LHS.getSInt(i) > (getSignedMaxValue(ResultWidth) - RHS.getSInt(i)))) || + ((RHS.getSInt(i) < 0) && (LHS.getSInt(i) < (getSignedMinValue(ResultWidth) - RHS.getSInt(i))))) + { + logArithmeticException(); + } + } + } + + break; + } + + //case llvm::Instruction::FAdd: + //case llvm::Instruction::FDiv: + //case llvm::Instruction::FMul: + + case llvm::Instruction::FPToSI: + { + const auto&& CastInst = llvm::dyn_cast(instruction); + const TypedValue Val = workItem->getOperand(instruction->getOperand(0)); + const unsigned int ResultWidth = CastInst->getDestTy()->getScalarSizeInBits(); + + if(Val.getFloat(0) < getSignedMinValue(ResultWidth) || Val.getFloat(0) > getSignedMaxValue(ResultWidth)) + { + logArithmeticException(); + } + + break; + } + case llvm::Instruction::FPToUI: + { + const auto&& CastInst = llvm::dyn_cast(instruction); + const TypedValue Val = workItem->getOperand(instruction->getOperand(0)); + const unsigned int ResultWidth = CastInst->getDestTy()->getScalarSizeInBits(); + + if(Val.getFloat(0) < 0 || Val.getFloat(0) > getUnsignedMaxValue(ResultWidth)) + { + logArithmeticException(); + } + + break; + } + case llvm::Instruction::FPTrunc: + { + const auto&& CastInst = llvm::dyn_cast(instruction); + const TypedValue Val = workItem->getOperand(instruction->getOperand(0)); + + if((CastInst->getDestTy()->isFloatTy() && abs(Val.getFloat(0)) > FLT_MAX) || + (CastInst->getDestTy()->isHalfTy() && abs(Val.getFloat(0)) > HALF_MAX)) + { + logArithmeticException(); + } + + break; + } + + //case llvm::Instruction::FRem: + //case llvm::Instruction::FSub: + + case llvm::Instruction::Mul: + { + const auto&& BinOp = llvm::dyn_cast(instruction); + + // Check for signed overflow + // Report an exception iff LLVM would create a poisoned value on overflow + // Other than that it is not possible to differentiate between signed/unsigned values + if(BinOp->hasNoSignedWrap()) + { + const TypedValue LHS = workItem->getOperand(instruction->getOperand(0)); + const TypedValue RHS = workItem->getOperand(instruction->getOperand(1)); + const unsigned int ResultWidth = BinOp->getType()->getScalarSizeInBits(); + + for(unsigned int i = 0; i < result.num; ++i) + { + if(LHS.getSInt(i) > 0) { /* LHS is positive */ + if(RHS.getSInt(i) > 0) { /* LHS and RHS are positive */ + if(LHS.getSInt(i) > (getSignedMaxValue(ResultWidth) / RHS.getSInt(i))) + { + logArithmeticException(); + } + } else { /* LHS positive, RHS nonpositive */ + if(RHS.getSInt(i) < (getSignedMinValue(ResultWidth) / LHS.getSInt(i))) + { + logArithmeticException(); + } + } + } else { /* LHS is nonpositive */ + if(RHS.getSInt(i) > 0) { /* LHS is nonpositive, RHS is positive */ + if(LHS.getSInt(i) < (getSignedMinValue(ResultWidth) / RHS.getSInt(i))) + { + logArithmeticException(); + } + } else { /* LHS and RHS are nonpositive */ + if((LHS.getSInt(i) != 0) && (RHS.getSInt(i) < (getSignedMaxValue(ResultWidth) / LHS.getSInt(i)))) + { + logArithmeticException(); + } + } + } + } + } + + break; + } + case llvm::Instruction::PtrToInt: + { + const auto&& PtrToIntInst = llvm::dyn_cast(instruction); + + // Check for signed overflow and division by zero + const TypedValue PtrValue = workItem->getOperand(PtrToIntInst->getPointerOperand()); + const unsigned int ResultWidth = PtrToIntInst->getDestTy()->getScalarSizeInBits(); + + for(unsigned int i = 0; i < result.num; ++i) + { + //TODO: Is this check ok? + if((uintptr_t)PtrValue.getPointer(i) > (uintptr_t)getSignedMaxValue(ResultWidth)) + { + logArithmeticException(); + } + } + + break; + } + case llvm::Instruction::SDiv: + { + const auto&& BinOp = llvm::dyn_cast(instruction); + + // Check for signed overflow and division by zero + const TypedValue LHS = workItem->getOperand(instruction->getOperand(0)); + const TypedValue RHS = workItem->getOperand(instruction->getOperand(1)); + const unsigned int ResultWidth = BinOp->getType()->getScalarSizeInBits(); + + for(unsigned int i = 0; i < result.num; ++i) + { + if((RHS.getSInt(i) == 0) || ((LHS.getSInt(i) == getSignedMinValue(ResultWidth)) && (RHS.getSInt(i) == -1))) + { + logArithmeticException(); + } + } + + break; + } + case llvm::Instruction::SIToFP: + { + const auto&& CastInst = llvm::dyn_cast(instruction); + const TypedValue Val = workItem->getOperand(instruction->getOperand(0)); + + double MaxValue = 0; + + if(CastInst->getDestTy()->isDoubleTy()) + { + MaxValue = DBL_MAX; + } + else if(CastInst->getDestTy()->isFloatTy()) + { + MaxValue = FLT_MAX; + } + else if(CastInst->getDestTy()->isHalfTy()) + { + MaxValue = HALF_MAX; + } + else + { + FATAL_ERROR("Unknown float type"); + } + + if(abs(Val.getSInt(0)) > MaxValue) + { + logArithmeticException(); + } + + break; + } + case llvm::Instruction::SRem: + { + const auto&& BinOp = llvm::dyn_cast(instruction); + + // Check for signed overflow and division by zero + const TypedValue LHS = workItem->getOperand(instruction->getOperand(0)); + const TypedValue RHS = workItem->getOperand(instruction->getOperand(1)); + const unsigned int ResultWidth = BinOp->getType()->getScalarSizeInBits(); + + for(unsigned int i = 0; i < result.num; ++i) + { + if((RHS.getSInt(i) == 0) || ((LHS.getSInt(i) == getSignedMinValue(ResultWidth)) && (RHS.getSInt(i) == -1))) + { + logArithmeticException(); + } + } + + break; + } + case llvm::Instruction::Sub: + { + const auto&& BinOp = llvm::dyn_cast(instruction); + + // Check for signed overflow + // Report an exception iff LLVM would create a poisoned value on overflow + // Other than that it is not possible to differentiate between signed/unsigned values + if(BinOp->hasNoSignedWrap()) + { + const TypedValue LHS = workItem->getOperand(instruction->getOperand(0)); + const TypedValue RHS = workItem->getOperand(instruction->getOperand(1)); + const unsigned int ResultWidth = BinOp->getType()->getScalarSizeInBits(); + + for(unsigned int i = 0; i < result.num; ++i) + { + + if((RHS.getSInt(i) > 0 && LHS.getSInt(i) < getSignedMinValue(ResultWidth) + RHS.getSInt(i)) || + (RHS.getSInt(i) < 0 && LHS.getSInt(i) > getSignedMaxValue(ResultWidth) + RHS.getSInt(i))) + { + logArithmeticException(); + } + } + } + + break; + } + case llvm::Instruction::UDiv: + { + // Check for division by zero + const TypedValue RHS = workItem->getOperand(instruction->getOperand(1)); + + for(unsigned int i = 0; i < result.num; ++i) + { + if(RHS.getSInt(i) == 0) + { + logArithmeticException(); + } + } + + break; + } + case llvm::Instruction::UIToFP: + { + const auto&& CastInst = llvm::dyn_cast(instruction); + const TypedValue Val = workItem->getOperand(instruction->getOperand(0)); + + double MaxValue = 0; + + if(CastInst->getDestTy()->isDoubleTy()) + { + MaxValue = DBL_MAX; + } + else if(CastInst->getDestTy()->isFloatTy()) + { + MaxValue = FLT_MAX; + } + else if(CastInst->getDestTy()->isHalfTy()) + { + MaxValue = HALF_MAX; + } + else + { + FATAL_ERROR("Unknown float type"); + } + + if(Val.getUInt(0) > MaxValue) + { + logArithmeticException(); + } + + break; + } + case llvm::Instruction::URem: + { + // Check for division by zero + const TypedValue RHS = workItem->getOperand(instruction->getOperand(1)); + + for(unsigned int i = 0; i < result.num; ++i) + { + if(RHS.getSInt(i) == 0) + { + logArithmeticException(); + } + } + + break; + } + case llvm::Instruction::Unreachable: + FATAL_ERROR("Encountered unreachable instruction"); + } +} + +void ArithmeticExceptions::logArithmeticException() const +{ + Context::Message msg(WARNING, m_context); + msg << "Undefined behaviour due to an arithmetic exception" << endl + << msg.INDENT + << "Kernel: " << msg.CURRENT_KERNEL << endl + << "Entity: " << msg.CURRENT_ENTITY << endl + << msg.CURRENT_LOCATION << endl; + msg.send(); +} + diff --git a/src/plugins/ArithmeticExceptions.h b/src/plugins/ArithmeticExceptions.h new file mode 100644 index 00000000..57285523 --- /dev/null +++ b/src/plugins/ArithmeticExceptions.h @@ -0,0 +1,27 @@ +// ArithmeticExceptions.h (Oclgrind) +// Copyright (c) 2015, Moritz Pflanzer, +// Imperial College London. All rights reserved. +// +// This program is provided under a three-clause BSD license. For full +// license terms please see the LICENSE file distributed with this +// source code. + +#include "core/Plugin.h" + +namespace oclgrind +{ + class ArithmeticExceptions : public Plugin + { + public: + ArithmeticExceptions(const Context *context) : Plugin(context){}; + + virtual void instructionExecuted(const WorkItem *workItem, + const llvm::Instruction *instruction, + const TypedValue& result) override; + //virtual void kernelBegin(const KernelInvocation *kernelInvocation) override; + //virtual void kernelEnd(const KernelInvocation *kernelInvocation) override; + + private: + void logArithmeticException() const; + }; +} diff --git a/src/runtime/oclgrind.cpp b/src/runtime/oclgrind.cpp index dfe5c7a2..94f8574c 100644 --- a/src/runtime/oclgrind.cpp +++ b/src/runtime/oclgrind.cpp @@ -198,7 +198,11 @@ static bool parseArguments(int argc, char *argv[]) { for (int i = 1; i < argc; i++) { - if (!strcmp(argv[i], "--build-options")) + if (!strcmp(argv[i], "--arithmetic-exceptions")) + { + setEnvironment("OCLGRIND_ARITHMETIC_EXCEPTIONS", "1"); + } + else if (!strcmp(argv[i], "--build-options")) { if (++i >= argc) { diff --git a/tests/kernels/CMakeLists.txt b/tests/kernels/CMakeLists.txt index 017acb25..caef46eb 100644 --- a/tests/kernels/CMakeLists.txt +++ b/tests/kernels/CMakeLists.txt @@ -24,4 +24,8 @@ set_tests_properties(${KERNEL_TESTS} PROPERTIES # Expected failures set_tests_properties( + arithmetic/add + arithmetic/mul + arithmetic/ptr + arithmetic/sub PROPERTIES WILL_FAIL TRUE) diff --git a/tests/kernels/TESTS b/tests/kernels/TESTS index fd952699..7a5c538b 100644 --- a/tests/kernels/TESTS +++ b/tests/kernels/TESTS @@ -1,5 +1,11 @@ alignment/packed alignment/unaligned +arithmetic/add +arithmetic/div +arithmetic/mul +arithmetic/ptr +arithmetic/rem +arithmetic/sub async_copy/async_copy async_copy/async_copy_divergent async_copy/async_copy_global_race diff --git a/tests/kernels/arithmetic/add.cl b/tests/kernels/arithmetic/add.cl new file mode 100644 index 00000000..dab70caf --- /dev/null +++ b/tests/kernels/arithmetic/add.cl @@ -0,0 +1,17 @@ +__kernel void add(__global uint *output) +{ + volatile int a = 3; + volatile int b = 5; + + volatile int c = 3; + volatile int d = INT_MAX; + + volatile int2 va = {3, 3}; + volatile int2 vb = {5, INT_MAX}; + + output[0] = a + b; + output[1] = c + d; + int2 t = va + vb; + output[2] = t.x; + output[3] = t.y; +} diff --git a/tests/kernels/arithmetic/add.ref b/tests/kernels/arithmetic/add.ref new file mode 100644 index 00000000..7c681148 --- /dev/null +++ b/tests/kernels/arithmetic/add.ref @@ -0,0 +1,8 @@ +ERROR Undefined behaviour due to an arithmetic exception +ERROR Undefined behaviour due to an arithmetic exception + +EXACT Argument 'output': 16 bytes +EXACT output[0] = 8 +MATCH output[1] = +EXACT output[2] = 8 +MATCH output[3] = diff --git a/tests/kernels/arithmetic/add.sim b/tests/kernels/arithmetic/add.sim new file mode 100644 index 00000000..5f21b057 --- /dev/null +++ b/tests/kernels/arithmetic/add.sim @@ -0,0 +1,6 @@ +add.cl +add +1 1 1 +1 1 1 + + diff --git a/tests/kernels/arithmetic/div.cl b/tests/kernels/arithmetic/div.cl new file mode 100644 index 00000000..d006a0f8 --- /dev/null +++ b/tests/kernels/arithmetic/div.cl @@ -0,0 +1,24 @@ +__kernel void div(__global uint *output) +{ + volatile int a = INT_MIN; + volatile int b = 2; + volatile int c = 0; + volatile int d = -1; + volatile uint e = 0; + + volatile int4 vb = {2, 0, -1, 1}; + volatile uint2 vc = {0, 1}; + + output[0] = a / b; + output[1] = a / c; + output[2] = a / d; + output[3] = a / e; + int4 t = a / vb; + output[4] = t.x; + output[5] = t.y; + output[6] = t.z; + output[7] = t.w; + uint2 u = a / vc; + output[8] = u.x; + output[9] = u.y; +} diff --git a/tests/kernels/arithmetic/div.ref b/tests/kernels/arithmetic/div.ref new file mode 100644 index 00000000..6b02318d --- /dev/null +++ b/tests/kernels/arithmetic/div.ref @@ -0,0 +1,18 @@ +ERROR Undefined behaviour due to an arithmetic exception +ERROR Undefined behaviour due to an arithmetic exception +ERROR Undefined behaviour due to an arithmetic exception +ERROR Undefined behaviour due to an arithmetic exception +ERROR Undefined behaviour due to an arithmetic exception +ERROR Undefined behaviour due to an arithmetic exception + +EXACT Argument 'output': 40 bytes +EXACT output[0] = -1073741824 +MATCH output[1] = +MATCH output[2] = +MATCH output[3] = +EXACT output[4] = -1073741824 +MATCH output[5] = +MATCH output[6] = +EXACT output[7] = -2147483648 +MATCH output[8] = +EXACT output[9] = -2147483648 diff --git a/tests/kernels/arithmetic/div.sim b/tests/kernels/arithmetic/div.sim new file mode 100644 index 00000000..00bcbeb1 --- /dev/null +++ b/tests/kernels/arithmetic/div.sim @@ -0,0 +1,6 @@ +div.cl +div +1 1 1 +1 1 1 + + diff --git a/tests/kernels/arithmetic/mul.cl b/tests/kernels/arithmetic/mul.cl new file mode 100644 index 00000000..34595dbe --- /dev/null +++ b/tests/kernels/arithmetic/mul.cl @@ -0,0 +1,16 @@ +__kernel void mul(__global uint *output) +{ + volatile int a = 3; + volatile int b = 5; + + volatile int c = 3; + volatile int d = INT_MAX; + + volatile int2 va = {3, INT_MAX}; + + output[0] = a * b; + output[1] = c * d; + int2 t = va * b; + output[2] = t.x; + output[3] = t.y; +} diff --git a/tests/kernels/arithmetic/mul.ref b/tests/kernels/arithmetic/mul.ref new file mode 100644 index 00000000..7c681148 --- /dev/null +++ b/tests/kernels/arithmetic/mul.ref @@ -0,0 +1,8 @@ +ERROR Undefined behaviour due to an arithmetic exception +ERROR Undefined behaviour due to an arithmetic exception + +EXACT Argument 'output': 16 bytes +EXACT output[0] = 8 +MATCH output[1] = +EXACT output[2] = 8 +MATCH output[3] = diff --git a/tests/kernels/arithmetic/mul.sim b/tests/kernels/arithmetic/mul.sim new file mode 100644 index 00000000..a4e6ee9c --- /dev/null +++ b/tests/kernels/arithmetic/mul.sim @@ -0,0 +1,6 @@ +mul.cl +mul +1 1 1 +1 1 1 + + diff --git a/tests/kernels/arithmetic/ptr.cl b/tests/kernels/arithmetic/ptr.cl new file mode 100644 index 00000000..bb2c992e --- /dev/null +++ b/tests/kernels/arithmetic/ptr.cl @@ -0,0 +1,14 @@ +__kernel void ptr(__global ushort *output) +{ + volatile ushort a; + volatile char b; + + volatile ushort2 va; + + a = &b; + va = (ushort2)(1, &b); + + output[0] = a; + output[1] = va.x; + output[2] = va.y; +} diff --git a/tests/kernels/arithmetic/ptr.ref b/tests/kernels/arithmetic/ptr.ref new file mode 100644 index 00000000..ad1063e0 --- /dev/null +++ b/tests/kernels/arithmetic/ptr.ref @@ -0,0 +1,7 @@ +ERROR Undefined behaviour due to an arithmetic exception +ERROR Undefined behaviour due to an arithmetic exception + +EXACT Argument 'output': 6 bytes +MATCH output[0] = +EXACT output[1] = 1 +MATCH output[2] = diff --git a/tests/kernels/arithmetic/ptr.sim b/tests/kernels/arithmetic/ptr.sim new file mode 100644 index 00000000..e867260c --- /dev/null +++ b/tests/kernels/arithmetic/ptr.sim @@ -0,0 +1,6 @@ +ptr.cl +ptr +1 1 1 +1 1 1 + + diff --git a/tests/kernels/arithmetic/rem.cl b/tests/kernels/arithmetic/rem.cl new file mode 100644 index 00000000..3348a00e --- /dev/null +++ b/tests/kernels/arithmetic/rem.cl @@ -0,0 +1,24 @@ +__kernel void rem(__global uint *output) +{ + volatile int a = INT_MIN; + volatile int b = 2; + volatile int c = 0; + volatile int d = -1; + volatile uint e = 0; + + volatile int4 vb = {2, 0, -1, 1}; + volatile uint2 vc = {0, 1}; + + output[0] = a % b; + output[1] = a % c; + output[2] = a % d; + output[3] = a % e; + int4 t = a % vb; + output[4] = t.x; + output[5] = t.y; + output[6] = t.z; + output[7] = t.w; + uint2 u = a % vc; + output[8] = u.x; + output[9] = u.y; +} diff --git a/tests/kernels/arithmetic/rem.ref b/tests/kernels/arithmetic/rem.ref new file mode 100644 index 00000000..95c16c8a --- /dev/null +++ b/tests/kernels/arithmetic/rem.ref @@ -0,0 +1,18 @@ +ERROR Undefined behaviour due to an arithmetic exception +ERROR Undefined behaviour due to an arithmetic exception +ERROR Undefined behaviour due to an arithmetic exception +ERROR Undefined behaviour due to an arithmetic exception +ERROR Undefined behaviour due to an arithmetic exception +ERROR Undefined behaviour due to an arithmetic exception + +EXACT Argument 'output': 40 bytes +EXACT output[0] = 0 +MATCH output[1] = +MATCH output[2] = +MATCH output[3] = +EXACT output[4] = 0 +MATCH output[5] = +MATCH output[6] = +EXACT output[7] = 0 +MATCH output[8] = +EXACT output[9] = 0 diff --git a/tests/kernels/arithmetic/rem.sim b/tests/kernels/arithmetic/rem.sim new file mode 100644 index 00000000..ca24a485 --- /dev/null +++ b/tests/kernels/arithmetic/rem.sim @@ -0,0 +1,6 @@ +rem.cl +rem +1 1 1 +1 1 1 + + diff --git a/tests/kernels/arithmetic/sub.cl b/tests/kernels/arithmetic/sub.cl new file mode 100644 index 00000000..a4a0fa06 --- /dev/null +++ b/tests/kernels/arithmetic/sub.cl @@ -0,0 +1,17 @@ +__kernel void sub(__global uint *output) +{ + volatile int a = 3; + volatile int b = 5; + + volatile int c = 3; + volatile int d = INT_MIN; + + volatile int2 va = {3, 3}; + volatile int2 vb = {5, INT_MIN}; + + output[0] = a - b; + output[1] = c - d; + int2 t = va - vb; + output[2] = t.x; + output[3] = t.y; +} diff --git a/tests/kernels/arithmetic/sub.ref b/tests/kernels/arithmetic/sub.ref new file mode 100644 index 00000000..10da25a6 --- /dev/null +++ b/tests/kernels/arithmetic/sub.ref @@ -0,0 +1,8 @@ +ERROR Undefined behaviour due to an arithmetic exception +ERROR Undefined behaviour due to an arithmetic exception + +EXACT Argument 'output': 16 bytes +EXACT output[0] = -2 +MATCH output[1] = +EXACT output[2] = -2 +MATCH output[3] = diff --git a/tests/kernels/arithmetic/sub.sim b/tests/kernels/arithmetic/sub.sim new file mode 100644 index 00000000..0b7f9f07 --- /dev/null +++ b/tests/kernels/arithmetic/sub.sim @@ -0,0 +1,6 @@ +sub.cl +sub +1 1 1 +1 1 1 + + diff --git a/tests/run_test.py b/tests/run_test.py index 330ca7e0..9abda645 100644 --- a/tests/run_test.py +++ b/tests/run_test.py @@ -43,6 +43,7 @@ os.environ["OCLGRIND_CHECK_API"] = "1" os.environ["OCLGRIND_DATA_RACES"] = "1" os.environ["OCLGRIND_UNINITIALIZED"] = "1" +os.environ["OCLGRIND_ARITHMETIC_EXCEPTIONS"] = "1" def fail(ret=1): print('FAILED')