[Bf-blender-cvs] [0226279d0e7] temp-cycles-denoising: Cycles Denoising: Fix AMD OpenCL support

[Lukas Stockner]

Commit: 0226279d0e72ec3d29a5853fac3a33883ee77986
Author: Lukas Stockner
Date:   Tue Apr 11 19:32:46 2017 +0200
Branches: temp-cycles-denoising
https://developer.blender.org/rB0226279d0e72ec3d29a5853fac3a33883ee77986

Cycles Denoising: Fix AMD OpenCL support

Credit for this patch goes to nirved.

===================================================================

M	intern/cycles/filter/CMakeLists.txt
M	intern/cycles/filter/filter_defines.h
M	intern/cycles/filter/filter_prefilter.h
M	intern/cycles/filter/kernels/opencl/filter.cl

===================================================================

diff --git a/intern/cycles/filter/CMakeLists.txt b/intern/cycles/filter/CMakeLists.txt
index 6d723583401..adea2cf1e26 100644
--- a/intern/cycles/filter/CMakeLists.txt
+++ b/intern/cycles/filter/CMakeLists.txt
@@ -170,6 +170,6 @@ endif()
 
 # OpenCL kernel
 
-delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/opencl/filter.cl" ${CYCLES_INSTALL_PATH}/kernel/kernels/opencl)
-delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/cuda/filter.cu" ${CYCLES_INSTALL_PATH}/kernel/kernels/cuda)
-delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_HEADERS}" ${CYCLES_INSTALL_PATH}/kernel)
+delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/opencl/filter.cl" ${CYCLES_INSTALL_PATH}/source/kernel/kernels/opencl)
+delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernels/cuda/filter.cu" ${CYCLES_INSTALL_PATH}/source/kernel/kernels/cuda)
+delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_HEADERS}" ${CYCLES_INSTALL_PATH}/source/filter)
diff --git a/intern/cycles/filter/filter_defines.h b/intern/cycles/filter/filter_defines.h
index e98596d21e8..4a55a91fc7f 100644
--- a/intern/cycles/filter/filter_defines.h
+++ b/intern/cycles/filter/filter_defines.h
@@ -29,7 +29,11 @@ typedef struct TilesInfo {
 	int x[4];
 	int y[4];
 	/* TODO(lukas): CUDA doesn't have uint64_t... */
+#ifdef __SPLIT_KERNEL__
+	ccl_global float *buffers[9];
+#else
 	long long int buffers[9];
+#endif
 } TilesInfo;
 
 #endif /* __FILTER_DEFINES_H__*/
diff --git a/intern/cycles/filter/filter_prefilter.h b/intern/cycles/filter/filter_prefilter.h
index e57600bde21..156992951d5 100644
--- a/intern/cycles/filter/filter_prefilter.h
+++ b/intern/cycles/filter/filter_prefilter.h
@@ -42,10 +42,10 @@ ccl_device void kernel_filter_divide_shadow(int sample,
 	int ytile = (y < tiles->y[1])? 0: ((y < tiles->y[2])? 1: 2);
 	int tile = ytile*3+xtile;
 
-	float ccl_readonly_ptr buffer = (float*) tiles->buffers[tile];
+	ccl_global float ccl_readonly_ptr buffer = (ccl_global float*) tiles->buffers[tile];
 	int offset = tiles->offsets[tile];
 	int stride = tiles->strides[tile];
-	float ccl_readonly_ptr center_buffer = buffer + (y*stride + x + offset)*buffer_pass_stride + buffer_denoising_offset;
+	ccl_global float ccl_readonly_ptr center_buffer = buffer + (y*stride + x + offset)*buffer_pass_stride + buffer_denoising_offset;
 
 	int buffer_w = align_up(rect.z - rect.x, 4);
 	int idx = (y-rect.y)*buffer_w + (x - rect.x);
@@ -90,7 +90,7 @@ ccl_device void kernel_filter_get_feature(int sample,
 	int xtile = (x < tiles->x[1])? 0: ((x < tiles->x[2])? 1: 2);
 	int ytile = (y < tiles->y[1])? 0: ((y < tiles->y[2])? 1: 2);
 	int tile = ytile*3+xtile;
-	float *center_buffer = ((float*) tiles->buffers[tile]) + (tiles->offsets[tile] + y*tiles->strides[tile] + x)*buffer_pass_stride + buffer_denoising_offset;
+	ccl_global float *center_buffer = ((ccl_global float*) tiles->buffers[tile]) + (tiles->offsets[tile] + y*tiles->strides[tile] + x)*buffer_pass_stride + buffer_denoising_offset;
 
 	int buffer_w = align_up(rect.z - rect.x, 4);
 	int idx = (y-rect.y)*buffer_w + (x - rect.x);
diff --git a/intern/cycles/filter/kernels/opencl/filter.cl b/intern/cycles/filter/kernels/opencl/filter.cl
index fd381f1e56f..e71c88fd54b 100644
--- a/intern/cycles/filter/kernels/opencl/filter.cl
+++ b/intern/cycles/filter/kernels/opencl/filter.cl
@@ -16,9 +16,9 @@
 
 /* OpenCL kernel entry points */
 
-#include "../../filter_compat_opencl.h"
+#include "filter/filter_compat_opencl.h"
 
-#include "../../filter_kernel.h"
+#include "filter/filter_kernel.h"
 
 /* kernels */
 
@@ -276,4 +276,4 @@ __kernel void kernel_ocl_filter_set_tiles(ccl_global TilesInfo* tiles,
 		tiles->buffers[7] = buffer_8;
 		tiles->buffers[8] = buffer_9;
 	}
-}
\ No newline at end of file
+}