From b83e16b7ffab072852d458d4a16661707ba026b6 Mon Sep 17 00:00:00 2001 From: "erwin.coumans" Date: Fri, 25 Jun 2010 00:40:20 +0000 Subject: [PATCH] fix out-of-bounds in AMD version of VectorAdd --- Demos/VectorAdd_OpenCL/MiniCL_VectorAdd.cpp | 3 ++- Demos/VectorAdd_OpenCL/VectorAddKernels.cl | 2 ++ src/LinearMath/btMinMax.h | 2 ++ 3 files changed, 6 insertions(+), 1 deletion(-) diff --git a/Demos/VectorAdd_OpenCL/MiniCL_VectorAdd.cpp b/Demos/VectorAdd_OpenCL/MiniCL_VectorAdd.cpp index 798eb7cc4..6220140e8 100644 --- a/Demos/VectorAdd_OpenCL/MiniCL_VectorAdd.cpp +++ b/Demos/VectorAdd_OpenCL/MiniCL_VectorAdd.cpp @@ -300,9 +300,10 @@ int main(int argc, char **argv) { num_t++; //this can cause problems -> processing outside of the buffer + //make sure to check kernel } - size_t globalThreads[] = {actualGlobalSize};//num_t * workgroupSize}; + size_t globalThreads[] = {num_t * workgroupSize}; size_t localThreads[] = {workgroupSize}; diff --git a/Demos/VectorAdd_OpenCL/VectorAddKernels.cl b/Demos/VectorAdd_OpenCL/VectorAddKernels.cl index a5031ec7e..70d96a93d 100644 --- a/Demos/VectorAdd_OpenCL/VectorAddKernels.cl +++ b/Demos/VectorAdd_OpenCL/VectorAddKernels.cl @@ -24,6 +24,8 @@ __kernel void VectorAdd(__global const float8* a, __global const float8* b, __gl { // get oct-float index into global data array int iGID = get_global_id(0); + if (iGID>=100000) + return; // read inputs into registers float8 f8InA = a[iGID]; diff --git a/src/LinearMath/btMinMax.h b/src/LinearMath/btMinMax.h index 5e27d62a4..33ea10ce3 100644 --- a/src/LinearMath/btMinMax.h +++ b/src/LinearMath/btMinMax.h @@ -17,6 +17,8 @@ subject to the following restrictions: #ifndef GEN_MINMAX_H #define GEN_MINMAX_H +#include "btScalar.h" + template SIMD_FORCE_INLINE const T& btMin(const T& a, const T& b) {