From 6d2c0af1db2586469eb9d1d1f07232dc2ae8fa8e Mon Sep 17 00:00:00 2001 From: Lu Guanqun Date: Wed, 10 Apr 2013 16:11:59 +0800 Subject: [PATCH] release the contraint of volatile pointer Signed-off-by: Lu Guanqun Reviewed-by: Zhigang Gong --- backend/src/llvm/llvm_gen_backend.cpp | 1 - kernels/compiler_volatile.cl | 4 ++++ utests/CMakeLists.txt | 1 + utests/compiler_volatile.cpp | 9 +++++++++ 4 files changed, 14 insertions(+), 1 deletion(-) create mode 100644 kernels/compiler_volatile.cl create mode 100644 utests/compiler_volatile.cpp diff --git a/backend/src/llvm/llvm_gen_backend.cpp b/backend/src/llvm/llvm_gen_backend.cpp index 6953aa8..42265ee 100644 --- a/backend/src/llvm/llvm_gen_backend.cpp +++ b/backend/src/llvm/llvm_gen_backend.cpp @@ -2093,7 +2093,6 @@ namespace gbe template INLINE void GenWriter::emitLoadOrStore(T &I) { - GBE_ASSERTM(I.isVolatile() == false, "Volatile pointer is not supported"); unsigned int llvmSpace = I.getPointerAddressSpace(); Value *llvmPtr = I.getPointerOperand(); Value *llvmValues = getLoadOrStoreValue(I); diff --git a/kernels/compiler_volatile.cl b/kernels/compiler_volatile.cl new file mode 100644 index 0000000..84f7228 --- /dev/null +++ b/kernels/compiler_volatile.cl @@ -0,0 +1,4 @@ +__kernel void compiler_volatile(__global int *dst, __local volatile int *hop) { + hop[get_global_id(0)] = get_local_id(1); + dst[get_global_id(0)] = hop[get_local_id(0)]; +} diff --git a/utests/CMakeLists.txt b/utests/CMakeLists.txt index a47dc24..efa1ae2 100644 --- a/utests/CMakeLists.txt +++ b/utests/CMakeLists.txt @@ -64,6 +64,7 @@ ADD_LIBRARY(utests SHARED compiler_local_memory_barrier.cpp compiler_local_memory_barrier_wg64.cpp compiler_movforphi_undef.cpp + compiler_volatile.cpp runtime_createcontext.cpp utest_assert.cpp utest.cpp diff --git a/utests/compiler_volatile.cpp b/utests/compiler_volatile.cpp new file mode 100644 index 0000000..f4fe054 --- /dev/null +++ b/utests/compiler_volatile.cpp @@ -0,0 +1,9 @@ +#include "utest_helper.hpp" + +void compiler_volatile(void) +{ + // Setup kernel and buffers + OCL_CREATE_KERNEL("compiler_volatile"); +} + +MAKE_UTEST_FROM_FUNCTION(compiler_volatile); -- 2.7.4