CS344 Intro To Parallel Programming Using CUDA lesson2

我们接下来介绍常见的Communication Patterns:

  1. Map: Tasks read from and write to specific data elements.

  1. Gather :

  1. Scatter: tasks compute where to write output

  2. Stencil: tasks read input from a fixed neighborhood in an array(Data reuse)

  3. Transpose:

接着我们对上述介绍到的Communication patterns 做下总结,包括接下来会涉猎到的reduce 和 scan/sort.

接下来我们将从硬件层面介绍,如何更高效的访问内存,以及进程之间如何通过共享内存进行通信。

从编程者的角度来看,GPU的编程模型如下:

其中我们将每个进程(thread)分配到不同的block中,这些block通过共同合作来解决一个子任务foo(); 同时GPU将每个thread blocks 分配到不同的SM上(程序里不用管),如下图所示:

我们可以确保,在foo() 完成之后,bar() 任务开始运行,但是不能保证foo()里thread block里的运行顺序(随机)。

总之,CUDA makes few guarantees about when and where thread blocks will run.

这样做的好处是:

  • hardware can run things efficiently

  • no waiting on slowpokes

  • scalability!

坏处有:

  • no assumptions blocks->SM
  • no communication between blocks ->”dead lock”
  • threads, blocks must complete.

CUDA guarantees that:

  • 在相同block里的所有进程在同一时间运行在相同的SM上。
  • 在下一个kernel运行之前,之前kernel里的所有block里的进程全部运行完成。

下面的例子可以说明这个问题:

#include <stdio.h>

#define NUM_BLOCKS 1
#define BLOCK_WIDTH 256

__global__ void hello()
{
    printf("Hello world! I'm thread %d\n", threadIdx.x);
}

int main(int argc,char **argv)
{
    // launch the kernel
    hello<<<NUM_BLOCKS, BLOCK_WIDTH>>>();

    // force the printf()s to flush
    cudaDeviceSynchronize();

    printf("That's all!\n");

    return 0;
}

我们可以发现,每次运行的结果都是随机的,这说明block 里 thread的完成先后顺序是随机的。

接下来我们介绍GPU的Memory Model.

我们可以看到,每个进程有自己独立的local memory, 也有每个Thread Block 共享的shared memory. 同时有所有Thread block 共享的 Global memory.

这也就意味着,每个thread可以通过shared memory 和 global memory 访问到其它进程的结果。这样会带来一定的风险:what if a thread reads a result before another thread writes it?

这也就说明了:threads needs to synchronize.

为了解决上述问题,我们引入Barrier: point in the program where threads stop and wait.

when all threads have reached the barrier, they can proceed.

这也就引入了下列的Programming Model:

我们通过下列代码来说明:

__global__ vodi foo(...) {
    // block里共享
    __shared__ int s[1024];
    int i = threadIdx.x;
    int tmp = s[i-1] + s[i] + s[i+1];
    __syncthreads();
    s[i] = tmp;
}

为了编写高效的程序,我们需要:

  • 最大化每个线程的操作数。

  • 最小化每个线程花在内存上的时间。

为了减少每个线程在内存上花费的时间,我们需要更快,更高效的访问内存:

  • local > shared » global » cpu host

另外,对于GPU来说,访问连续的内存更高效,如下图所示:

另外,我们考虑一个新问题:如果我们用10000个线程来增加长度为10的数组,每次加一,该如何实现?

我们的常规做法如下:

__global__ void increment_naive(int *g)
{
	// which thread is this?
	int i = blockIdx.x * blockDim.x + threadIdx.x; 

	// each thread to increment consecutive elements, wrapping at ARRAY_SIZE
	i = i % ARRAY_SIZE;  
	g[i] = g[i] + 1;
}

这样做不会得到我们想要的结果,因为这样会造成 访问冲突。因此我们需要使用atomicAdd函数:

__global__ void increment_atomic(int *g)
{
	// which thread is this?
	int i = blockIdx.x * blockDim.x + threadIdx.x; 

	// each thread to increment consecutive elements, wrapping at ARRAY_SIZE
	i = i % ARRAY_SIZE;  
	atomicAdd(& g[i], 1);
}

完整的实现代码如下:

#include <stdio.h>
#include "gputimer.h"

#define NUM_THREADS 1000000
#define ARRAY_SIZE  100

#define BLOCK_WIDTH 1000

void print_array(int *array, int size)
{
    printf("{ ");
    for (int i = 0; i < size; i++)  { printf("%d ", array[i]); }
    printf("}\n");
}

__global__ void increment_atomic(int *g)
{
	// which thread is this?
	int i = blockIdx.x * blockDim.x + threadIdx.x; 

	// each thread to increment consecutive elements, wrapping at ARRAY_SIZE
	i = i % ARRAY_SIZE;  
	atomicAdd(& g[i], 1);
}

int main(int argc,char **argv)
{   
    GpuTimer timer;
    printf("%d total threads in %d blocks writing into %d array elements\n",
           NUM_THREADS, NUM_THREADS / BLOCK_WIDTH, ARRAY_SIZE);

    // declare and allocate host memory
    int h_array[ARRAY_SIZE];
    const int ARRAY_BYTES = ARRAY_SIZE * sizeof(int);
 
    // declare, allocate, and zero out GPU memory
    int * d_array;
    cudaMalloc((void **) &d_array, ARRAY_BYTES);
    cudaMemset((void *) d_array, 0, ARRAY_BYTES); 

    // launch the kernel - comment out one of these
    timer.Start();
    // increment_naive<<<NUM_THREADS/BLOCK_WIDTH, BLOCK_WIDTH>>>(d_array);
    increment_atomic<<<NUM_THREADS/BLOCK_WIDTH, BLOCK_WIDTH>>>(d_array);
    timer.Stop();
    
    // copy back the array of sums from GPU and print
    cudaMemcpy(h_array, d_array, ARRAY_BYTES, cudaMemcpyDeviceToHost);
    print_array(h_array, ARRAY_SIZE);
    printf("Time elapsed = %g ms\n", timer.Elapsed());
 
    // free GPU memory allocation and exit
    cudaFree(d_array);
    return 0;
}

此外,atomic函数也有缺点,具体如下:

  • 只支持特定函数和输入类型
  • 对于float类型,依然没有顺序不变性 ($(a+b)+c != a + (b+c)$)
  • 由于是串行操作,因此速度会较慢。

同时,为了编写高效的程序,我们需要避免线程发散,具体如下:

总结: