document the abstract execution interface

Author: SimonDanisch

docs/src/index.md

@@ -1,5 +1,46 @@
 # GPUArrays Documentation
 
+
+# Abstract GPU interface
+
+GPUArrays supports different platforms like CUDA and OpenCL, which all have different
+names for function that offer the same functionality on the hardware.
+E.g. how to call a function on the GPU, how to get the thread index etc.
+GPUArrays offers an abstract interface for these functions which are overloaded
+by the packages like [CLArrays](https://github.com/JuliaGPU/CLArrays.jl) and [CuArrays](https://github.com/JuliaGPU/CuArrays.jl).
+This makes it possible to write generic code that can be run on all hardware.
+GPUArrays itself even contains a pure Julia implementation of this interface.
+The julia reference implementation is also a great way to debug your GPU code, since it
+offers many more errors and debugging information compared to the GPU backends - which
+mostly silently error or give cryptic errors (so far).
+You can use the reference implementation by using the `GPUArrays.JLArray` type.
+
+The functions that are currently part of the interface:
+
+The low level dim + idx function, with a similar naming as in CUDA (with `*` indicating `(x, y, z)`):
+```Julia
+blockidx_*(state), blockdim_*(state), threadidx_*(state), griddim_*(state)
+# Known in OpenCL as:
+get_group_id,      get_local_size,    get_local_id,       get_num_groups
+```
+
 ```@docs
 gpu_call(f, A::GPUArray, args::Tuple, configuration = length(A))
+
+
+linear_index(state)
+
+global_size(state)
+
+@linearidx(A, statesym = :state)
+
+@cartesianidx(A, statesym = :state)
+
+
+synchronize_threads(state)
+
+
+device(A::AbstractArray)
+synchronize(A::AbstractArray)
+
 ```

src/abstract_gpu_interface.jl

@@ -25,45 +25,60 @@ linear index in a GPU kernel (equal to  OpenCL.get_global_id)
 @inline function linear_index(state)
     UInt32((blockidx_x(state) - UInt32(1)) * blockdim_x(state) + threadidx_x(state))
 end
+
+"""
+Macro form of `linear_index`, which returns when out of bounds
+"""
+macro linearidx(A, statesym = :state)
+    quote
+        x1 = $(esc(A))
+        i1 = linear_index($(esc(statesym)))
+        i1 > length(x1) && return
+        i1
+    end
+end
+
+
+"""
+Like `@linearidx`, but returns an N-dimensional `NTuple{ndim(A), Cuint}` as index
+"""
+macro cartesianidx(A, statesym = :state)
+    quote
+        x = $(esc(A))
+        i2 = @linearidx(x, $(esc(statesym)))
+        gpu_ind2sub(x, i2)
+    end
+end
+
+"""
+Global size == blockdim * griddim == total number of kernel execution
+"""
 @inline function global_size(state)
+    # TODO nd version
     griddim_x(state) * blockdim_x(state)
 end
 
+
 """
-Blocks until all operations are finished on `A`
+Gets the device associated to the Array `A`
 """
-function synchronize(A::AbstractArray)
+function device(A::AbstractArray)
     # fallback is a noop, for backends not needing synchronization. This
     # makes it easier to write generic code that also works for AbstractArrays
 end
 """
-Gets the device associated to the Array `A`
+Blocks until all operations are finished on `A`
 """
-function device(A::AbstractArray)
+function synchronize(A::AbstractArray)
     # fallback is a noop, for backends not needing synchronization. This
     # makes it easier to write generic code that also works for AbstractArrays
 end
-
 #
 # @inline function synchronize_threads(state)
 #     CUDAnative.__syncthreads()
 # end
 
-macro linearidx(A, statesym = :state)
-    quote
-        x1 = $(esc(A))
-        i1 = linear_index($(esc(statesym)))
-        i1 > length(x1) && return
-        i1
-    end
-end
-macro cartesianidx(A, statesym = :state)
-    quote
-        x = $(esc(A))
-        i2 = @linearidx(x, $(esc(statesym)))
-        gpu_ind2sub(x, i2)
-    end
-end
+
 
 
 """