My experience is in porting numerical algorithms to NVIDIA GPUs with the CUDA platform. If your intention is to decide whether a GPU port is worth it, then perh

My experience is in porting numerical algorithms to NVIDIA GPUs with the CUDA platform. If your intention is to decide whether a GPU port is worth it, then perhaps you should be thinking at a higher level of abstraction. I only "speak" CUDA, however the primitives provided by the CUDA/OpenCL development platforms are largely equivalent. For a single "kernel" the GPU provides parallelism at two levels, 1. Fine-grained Thread-level parallelism: You have a group of threads performing ideally identical operations. The model is called SIMT (Single Instruction Multiple Thread) because it is not quite SIMD but close. Say you have (1-1024) threads, the hardware splits these into warps of 32 threads. The 32 threads within a warp execute together in classic SIMD manner, however only one of these warps is executing on the actual hardware at any given time. When a warp makes a request for data from high-latency device memory, it is typically suspended and another warp swapped in. This allows latency hiding between warps. However, calling __syncthreads() forces all warps to synchronise. Any conditional statements cause warp divergence, the warp is split and the alternative execution paths run sequentially. Clearly, this incurs a performance penalty as SIMD behaviour is lost and should be avoided as much as possible. The 'ideal' identical operation size is therefore a warp or 32 consecutive threads, however occasional divergence is acceptable. Memory access patterns are much more important however. 2. Coarse-grained Block-level parallelism. A block is an independent group of threads. The platform allows mutiple blocks to execute on different Streaming Multiprocessors. The hardware may schedule multiple blocks to execute on a single SM if there are sufficient resources (such as register and shared memory), this also facilitates latency hiding. Other than the kernel invocation which created multiple threads blocks, they are unrelated and cannot communicate or be synchronised. The potential thread level parallelism is what should be considered. If your algorithm is representable as fine-grained floating point vector operations without excessive divergence then it is well-suited to the GPU. Of course, the actual performance gain depends on the quality of implementation and optimisation is highly GPU-specific and non-trivial. I recommend http://developer.download.nvidia.com/compute/DevZone/C/html/C/src/reduction/doc/reduction.pdf http://developer.download.nvidia.com/compute/DevZone/C/html/... as an example of how important optimisations lead to those impressive performance numbers.