diff --git a/RTCP/Cobalt/GPUProc/CMakeLists.txt b/RTCP/Cobalt/GPUProc/CMakeLists.txt
index 219c987e0bf3c8d07f2f5897c693673e59445f50..b5958b5382d44f48755d4f6e5225bfc2f8464570 100644
--- a/RTCP/Cobalt/GPUProc/CMakeLists.txt
+++ b/RTCP/Cobalt/GPUProc/CMakeLists.txt
@@ -71,5 +71,5 @@ endif()
 
 add_subdirectory(src)
 add_subdirectory(test)
-add_subdirectory(share/gpu/kernels)
+add_subdirectory(share)
 add_subdirectory(etc)
diff --git a/RTCP/Cobalt/GPUProc/share/CMakeLists.txt b/RTCP/Cobalt/GPUProc/share/CMakeLists.txt
new file mode 100644
index 0000000000000000000000000000000000000000..3c173bbfe11b236f4128e014286da3e67e2a23aa
--- /dev/null
+++ b/RTCP/Cobalt/GPUProc/share/CMakeLists.txt
@@ -0,0 +1,6 @@
+# $Id$
+
+# Detect list of kernel sources to install.
+# CMake cannot detect changes in this list, but for kernel sources it doesn't need to.
+file(GLOB _kernel_sources RELATIVE "${CMAKE_CURRENT_SOURCE_DIR}" "gpu/kernels/*.cu" "gpu/kernels/*.cuh" "gpu/kernels/*.cl")
+lofar_add_data_files(${_kernel_sources})
diff --git a/RTCP/Cobalt/GPUProc/share/gpu/kernels/CMakeLists.txt b/RTCP/Cobalt/GPUProc/share/gpu/kernels/CMakeLists.txt
deleted file mode 100644
index 3a3a25b2f1863fe02046d05d86a91174871cf38b..0000000000000000000000000000000000000000
--- a/RTCP/Cobalt/GPUProc/share/gpu/kernels/CMakeLists.txt
+++ /dev/null
@@ -1,13 +0,0 @@
-# $Id$
-
-# Detect list of kernel sources to install.
-# CMake cannot detect changes in this list, but for kernel sources it doesn't need to.
-file(GLOB _kernel_path_sources "*.cu" "*.cuh" "*.cl")
-set(_kernel_sources )
-foreach(path ${_kernel_path_sources})
-  get_filename_component(filename "${path}" NAME)  # basename() (optional)
-  list(APPEND _kernel_sources "${filename}")
-endforeach()
-
-install(FILES ${_kernel_sources} DESTINATION share/gpu/kernels)
-
diff --git a/RTCP/Cobalt/GPUProc/share/gpu/kernels/Correlator.cu b/RTCP/Cobalt/GPUProc/share/gpu/kernels/Correlator.cu
index 8ce03145e646d9ca9763fb6e964a6b681d8fdb93..6bcb6ff231869bd5a1ce3bb5bf59e1bb0bbcaa11 100644
--- a/RTCP/Cobalt/GPUProc/share/gpu/kernels/Correlator.cu
+++ b/RTCP/Cobalt/GPUProc/share/gpu/kernels/Correlator.cu
@@ -27,6 +27,7 @@
  */
 
 #include "gpu_math.cuh"
+#include <assert.h>
 
 #define NR_BASELINES     (NR_STATIONS * (NR_STATIONS + 1) / 2)
 
@@ -198,6 +199,9 @@ __device__ void correlate_1x1(void *visibilitiesPtr, const void *correctedDataPt
   int x = ::major(baseline);
   int y = ::minor(baseline, x);
 
+  assert(x >= y);
+  assert(::baseline(x,y) == baseline);
+
   /* NOTE: stat0 >= statA holds */
   int stat_0 = x;
   int stat_A = y;
@@ -292,6 +296,9 @@ __device__ void correlate_2x2(void *visibilitiesPtr, const void *correctedDataPt
   int x = ::major(block);
   int y = ::minor(block, x);
 
+  assert(x >= y);
+  assert(::baseline(x,y) == block);
+
   /* NOTE: stat_0 >= stat_A holds */
   int stat_0 = 2 * x;
   int stat_A = 2 * y;
@@ -414,6 +421,9 @@ __device__ void correlate_3x3(void *visibilitiesPtr, const void *correctedDataPt
   int x = ::major(block);
   int y = ::minor(block, x);
 
+  assert(x >= y);
+  assert(::baseline(x,y) == block);
+
   /* NOTE: stat_0 >= stat_A holds */
   int stat_0 = 3 * x;
   int stat_A = 3 * y;
@@ -612,6 +622,9 @@ __device__ void correlate_4x4(void *visibilitiesPtr, const void *correctedDataPt
   int x = ::major(block);
   int y = ::minor(block, x);
 
+  assert(x >= y);
+  assert(::baseline(x,y) == block);
+
   /* NOTE: stat_0 >= stat_A holds */
   int stat_0 = 4 * x;
   int stat_A = 4 * y;