release the contraint of volatile pointer

Signed-off-by: Lu Guanqun <guanqun.lu@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>

diff --git a/backend/src/llvm/llvm_gen_backend.cpp b/backend/src/llvm/llvm_gen_backend.cpp
index 6953aa8..42265ee 100644 (file)
--- a/backend/src/llvm/llvm_gen_backend.cpp
+++ b/backend/src/llvm/llvm_gen_backend.cpp
@@ -2093,7 +2093,6 @@ namespace gbe
   template <bool isLoad, typename T>
   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 (file)
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 (file)
--- 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 (file)
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);