CUDA-5

CUDA-5

这次主要关注的原子操作-Atomic operations.

先看一下这个运算x++; 实际上这个操作经历了下面的三个过程

  • 读x的值

  • 在x的值上面加1

  • 把加完后的值从新写入x

这个过程可以称为 read-modify-write. 上面的过程在single-thread的时候是没有问题的,但是在multi-thread将会有问题,就是说同一个时刻可能有多个线程都在修改x的值,这将会变得出乎意料,接下来用一个例子来说明原子操作的心要性。 这个例子是直方图的计算,先看看在cpu上情况 这个文件的名字是cal_histogram.c.

#include <stdio.h>
#include <time.h>
#include <stdlib.h>

#define SIZE (100*1024*1024)

//generate random function

void* big_random_block(int size){
    unsigned char *data = (unsigned char*)malloc(size);
    for(int i=0; i<size; i++){
        data[i] = rand();
    }   
    return data;
}


int main(){
    unsigned char *buffer = (unsigned char*) big_random_block(SIZE);

    clock_t start, stop;
    start = clock();
    
    unsigned int histo[256];
    for(int i=0;i<256;i++){
        histo[i] = 0;
    }   
    
    for(int i=0;i<SIZE; i++)
        histo[buffer[i]]++;

    stop = clock();
    float elapsedTime = (float)(stop-start)/(float)CLOCKS_PER_SEC*1000.0f;

    printf("Time to generate: %3.1f ms\n", elapsedTime);

    //check the result
    long histoCount = 0;
    for(int i=0;i<256;i++)
        histoCount += histo[i];

    printf("Histogram sum: %ld\n", histoCount);
    free(buffer);

    return 0;


}

其中如果不加stdlib.h free那里会报错,不加time.h CLOCKS_PER_SEC那里会报错。

也可以不include上面的两个库,而把文件名改成.cu用nvcc进行编译也是可以通过的。

运行结果如下,

Time to generate: 150.1 ms
Histogram sum: 104857600

也可以看一下产生的数字

void* big_random_block(int size){
    unsigned char *data = (unsigned char*)malloc(size);
    for(int i=0; i<size; i++){
        data[i] = rand();
        printf("%d\n", data[i]);
    }   
    return data;
}


256 是因为each random 8-bit byte can be any of 256 different values(from 0x00 to 0xFF),所以需要256个桶来统计每个数字出现的次数。

接下来是gpu的version了,

这时候就会出现同一时刻可能多个线程都想修改同一个桶内的数字。

代码如下

#include <stdio.h>
#include <stdlib.h>

#define SIZE (100*1024*1024)

//generate random function

void* big_random_block(int size){
    unsigned char *data = (unsigned char*)malloc(size);
    for(int i=0; i<size; i++){
        data[i] = rand();
    }   
    return data;
}


// our kernel function

__global__ void histo_kernel(unsigned char *buffer, long size, unsigned int *histo){

    int i = threadIdx.x + blockIdx.x*blockDim.x;
    int stride = blockDim.x*gridDim.x;
    while(i<size){
        atomicAdd(&histo[buffer[i]], 1); 
        i += stride;
    }   
}

int main(void){
    unsigned char *buffer = (unsigned char*) big_random_block(SIZE);

    clock_t start, stop;
    start = clock();
    
    unsigned char *dev_buffer;
    unsigned int *dev_histo;
    
    //allocate 
    cudaMalloc((void**)&dev_buffer, SIZE);
    cudaMalloc((void**)&dev_histo, 256*sizeof(int));

    //cp data from cpu to gpu
    cudaMemcpy(dev_buffer, buffer, SIZE, cudaMemcpyHostToDevice);
    
    cudaMemset(dev_histo, 0, 256*sizeof(int));
    
    cudaDeviceProp prop;
    cudaGetDeviceProperties(&prop, 0); 
    int blocks = prop.multiProcessorCount;
    //printf("%d\n", blocks);
    histo_kernel<<<blocks*2, 256>>>(dev_buffer, SIZE, dev_histo);   // 2倍的mps数量最好。
    
    //get the result
    unsigned int histo[256];
    cudaMemcpy(histo, dev_histo, 256*sizeof(int), cudaMemcpyDeviceToHost);
    
    
    stop = clock();
    float elapsedTime = (float)(stop-start)/(float)CLOCKS_PER_SEC*1000.0f;
    printf("Time to generate: %3.1f ms\n", elapsedTime);


    //check the result
    long histoCount = 0;
    for(int i=0;i<256;i++)
        histoCount += histo[i];

    printf("Histogram sum: %ld\n", histoCount);
    
    cudaFree(dev_histo);
    cudaFree(dev_buffer);
    free(buffer);

    return 0;


}

运行结果如下,

Time to generate: 178.0 ms
Histogram sum: 104857600

先不讲结果对不对,但是从时间上来看竟然比CPU的还要慢,这是不能够容忍的事情,问题出在了哪里呢? 回头看一下核函数,由于是原子操作,那么当其中一个线程在改变值的时候,其他的线程是无法进行操作的,那这样就相当于和single-thread差不多了,何心搞那么复杂的。 实际上可以用上次里面的shared memory来处理这个问题。相当于是把任务分给每个blocks去做,这样不同的blocks在做的时候就不用管其它的blocks如何操作的,只要完成任务就行,而刚才的写法,其他的blocks里的线程都要等这个blocks里的某个线程写完之后才能够运行,结果很多的时间都花在了等待上面去了,而现在是不同的blocks中的线程不需要相互等待。 代码如下。

__global__ void histo_kernel( unsigned char *buffer,
                              long size,
                              unsigned int *histo ) { 

    __shared__  unsigned int temp[256];
    temp[threadIdx.x] = 0;
    __syncthreads();

    int i = threadIdx.x + blockIdx.x * blockDim.x;
    int stride = blockDim.x * gridDim.x;
    while (i < size) {
        atomicAdd( &temp[buffer[i]], 1 );
        i += stride;
    }   
    __syncthreads();
    atomicAdd( &(histo[threadIdx.x]), temp[threadIdx.x] );  // 汇总该block中的结果。
}



其他的地方不用变。

结果如下,颤抖吧,

Time to generate:  26.4 ms
Histogram Sum:  104857600

打赏,谢谢~~

取消

感谢您的支持,我会继续努力的!

扫码支持
扫码打赏,多谢支持~

打开微信扫一扫,即可进行扫码打赏哦