GPU 硬件有两个特殊的优势：原始计算 (FLOP) 和内存带宽。最困难的计算问题属于这两类之一。例如，稠密线性代数（A * B = C 或 Solve[Ax = y] 或 Diagonalize[A] 等）根据系统大小落在计算/内存带宽频谱的某个位置。快速傅里叶变换 (FFT) 也适合这种具有高聚合带宽需求的模型。与其他转换、基于网格/网格的算法、蒙特卡洛等一样。如果您查看 NVIDIA SDK代码示例，您可以了解最常解决的各种问题。

我认为更有启发性的答案是对“GPU 真正不擅长什么类型的问题？”这个问题。大多数不属于这一类的问题都可以在 GPU 上运行，尽管有些问题比其他问题需要更多的努力。

映射不好的问题通常太小或太不可预测。非常小的问题缺乏使用 GPU 上所有线程所需的并行性和/或可能适合 CPU 上的低级缓存，从而大大提高了 CPU 性能。不可预测的问题有太多有意义的分支，这可能会阻止数据从 GPU 内存有效地流式传输到内核，或者通过打破SIMD范式来降低并行度（参见“发散扭曲”）。此类问题的示例包括：

• 大多数图形算法（太不可预测，尤其是在内存空间中）
• 稀疏线性代数（但这对 CPU 也不利）
• 小信号处理问题（例如小于 1000 点的 FFT）
• 搜索
• 种类

具有高算术强度和常规内存访问模式的问题通常很容易（更）在 GPU 上实现，并且在它们上表现良好。

拥有高性能 GPU 代码的基本困难是您拥有大量内核，并且您希望它们都尽可能地发挥其全部效力。具有不规则内存访问模式或算术强度不高的问题使这变得困难：要么花费很长时间传达结果，要么花费很长时间从内存中获取内容（这很慢！），并且没有足够的时间处理数字。当然，代码中的并发潜力对于它在 GPU 上的良好实现能力也至关重要。

这本身并不是一个答案，而是对maxhutch和Reid.Atcheson其他答案的补充。

为了充分利用 GPU，您的问题不仅需要高度（或大规模）并行，而且将在 GPU 上执行的核心算法也应该尽可能小。在OpenCL术语中，这主要被称为内核。

更准确地说，内核应该适合GPU的每个多处理单元（或计算单元）的寄存器。寄存器的确切大小取决于 GPU。

鉴于内核足够小，问题的原始数据需要适合 GPU 的本地内存（读取：计算单元的本地内存(OpenCL) 或共享内存(CUDA)）。否则，即使 GPU 的高内存带宽也不够快，无法让处理元件一直处于忙碌状态。
通常这个内存大约 16 到 32 KiByte大。

可能是对先前回复的技术性补充：CUDA（即 Nvidia）GPU 可以描述为一组处理器，每个处理器在 32 个线程上自主工作。每个处理器中的线程以锁步方式工作（想想向量长度为​​ 32 的 SIMD）。

尽管使用 GPU 最诱人的方式是假装一切都在同步运行，但这并不总是最有效的做事方式。

如果您的代码不能很好地/自动地并行化到成百上千个线程，您可以将其分解为能够很好地并行化的单独异步任务，并执行那些只有 32 个线程以锁步方式运行的任务。CUDA 提供了一组原子指令，使实现互斥锁成为可能，而互斥锁又允许处理器在它们之间进行同步并在线程池范例中处理任务列表。然后，您的代码将以与多核系统上相同的方式工作，只需记住每个内核都有自己的 32 个线程。

这是一个使用 CUDA 的小例子，说明它是如何工作的

/* Global index of the next available task, assume this has been set to
   zero before spawning the kernel. */
__device__ int next_task;

/* We will use this value as our mutex variable. Assume it has been set to
   zero before spawning the kernel. */
__device__ int tasks_mutex;

/* Mutex routines using atomic compare-and-set. */
__device__ inline void cuda_mutex_lock ( int *m ) {
    while ( atomicCAS( m , 0 , 1 ) != 0 );
    }
__device__ inline void cuda_mutex_unlock ( int *m ) {
    atomicExch( m , 0 );
    }

__device__ void task_do ( struct task *t ) {

    /* Do whatever needs to be done for the task t using the 32 threads of
       a single warp. */
    }

__global__ void main ( struct task *tasks , int nr_tasks ) {

    __shared__ task_id;

    /* Main task loop... */
    while ( next_task < nr_tasks ) {

        /* The first thread in this block is responsible for picking-up a task. */
        if ( threadIdx.x == 0 ) {

            /* Get a hold of the task mutex. */
            cuda_mutex_lock( &tasks_mutex );

            /* Store the next task in the shared task_id variable so that all
               threads in this warp can see it. */
            task_id = next_task;

            /* Increase the task counter. */
            next_tast += 1;

            /* Make sure those last two writes to local and global memory can
               be seen by everybody. */
            __threadfence();

            /* Unlock the task mutex. */
            cuda_mutex_unlock( &tasks_mutex );

            }

        /* As of here, all threads in this warp are back in sync, so if we
           got a valid task, perform it. */
        if ( task_id < nr_tasks )
            task_do( &tasks[ task_id ] );

        } /* main loop. */

    }

然后，您必须调用内核main<<<N,32>>>(tasks,nr_tasks)以确保每个块仅包含 32 个线程，因此适合单个 warp。在这个例子中，为了简单起见，我还假设任务没有任何依赖关系（例如，一个任务依赖于另一个任务的结果）或冲突（例如，在同一个全局内存上工作）。如果是这样的话，那么任务选择就变得有点复杂了，但是结构本质上是一样的。

当然，这比只在一大批单元上做所有事情要复杂得多，但显着拓宽了 GPU 可以使用的问题类型。