[gtksourceview] cuda.lang: test the kernel-call context

[Sébastien Wilmet <swilmet src gnome org>]

commit 3d3ccbcc3e284138046c8089bbe4eb63a38ecfd4
Author: Sébastien Wilmet <swilmet gnome org>
Date:   Wed Nov 5 20:56:07 2014 +0100

    cuda.lang: test the kernel-call context
    
    When parsing the cuda.lang file there is a warning about the
    style-ref="latex:display-math", but it actually works (with a style
    scheme that defines a style for latex:display-math, such as tango and
    classic).
    
    The code is taken from (it is not the complete function):
    http://docs.nvidia.com/cuda/cuda-samples/#cuda-parallel-reduction
    
    https://bugzilla.gnome.org/show_bug.cgi?id=739684

 tests/syntax-highlighting/file.cu |   26 ++++++++++++++++++++++++++
 1 files changed, 26 insertions(+), 0 deletions(-)
---
diff --git a/tests/syntax-highlighting/file.cu b/tests/syntax-highlighting/file.cu
index dfe341a..630b0c4 100644
--- a/tests/syntax-highlighting/file.cu
+++ b/tests/syntax-highlighting/file.cu
@@ -12,3 +12,29 @@ void convertToArray(int **matrix, int *array, int N){
                for(unsigned int j=0; j< N; j++)
                        array[i*N+ j] = matrix[i][j];
 }
+
+// Wrapper function for kernel launch (not the complete function, just an example).
+template <class T>
+void
+reduce(int size, int threads, int blocks,
+       int whichKernel, T *d_idata, T *d_odata)
+{
+       dim3 dimBlock(threads, 1, 1);
+       dim3 dimGrid(blocks, 1, 1);
+
+       // when there is only one warp per block, we need to allocate two warps
+       // worth of shared memory so that we don't index shared memory out of bounds
+       int smemSize = (threads <= 32) ? 2 * threads * sizeof(T) : threads * sizeof(T);
+
+       // choose which of the optimized versions of reduction to launch
+       switch (whichKernel)
+       {
+               case 0:
+                       reduce0<T><<< dimGrid, dimBlock, smemSize >>>(d_idata, d_odata, size);
+                       break;
+
+               case 1:
+                       reduce1<T><<< dimGrid, dimBlock, smemSize >>>(d_idata, d_odata, size);
+                       break;
+       }
+}