OpenACC 简单的直方图

▶ 简单的直方图，强调原子操作的使用

● 代码

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

int main()
{
    const int length = 1024, basket = 10;
    int data[length], histgram[basket];

    srand(97);
    for (int i = 0; i < basket; histgram[i++] = 0);
    for (int i = 0; i < length; data[i++] = rand() % 10);

#pragma acc parallel loop
    for (int i = 0; i < length; i++)
    {
#pragma acc atomic update
        histgram[data[i]]+=1;
    }

    for (int i = 0; i < basket; i++)
        printf("histgram[%d] = %d
", i, histgram[i]);

    getchar();
    return 0;
}

● 输出结果，在 Windows 里是错的，在 WSL 里是对的

D:CodeOpenACCOpenACCProjectOpenACCProject>pgcc -acc -Minfo main.c -o main_acc.exe
main:
     14, Accelerator kernel generated
         Generating Tesla code
         15, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
     14, Generating implicit copyout(histgram[data])
         Generating implicit copyin(data[:1024])

D:CodeOpenACCOpenACCProjectOpenACCProject>main_acc.exe
launch CUDA kernel  file=D:CodeOpenACCOpenACCProjectOpenACCProjectmain.c function=main line=14 device=0 threadid=1 num_gangs=8 num_workers=1 vector_length=128 grid=8 block=128
histgram[0] = 0
histgram[1] = 0
histgram[2] = 0
histgram[3] = 0
histgram[4] = 0
histgram[5] = 100
histgram[6] = 0
histgram[7] = 0
histgram[8] = 0
histgram[9] = 0

cuan@CUAN:/mnt/d/Code/OpenACC/OpenACCProject/OpenACCProject$ pgcc -acc -Minfo main.c -o main_acc_ubuntu.exe
main:
     14, Accelerator kernel generated
         Generating Tesla code
         15, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
     14, Generating implicit copyout(histgram[data])
         Generating implicit copyin(data[:1024])
cuan@CUAN:/mnt/d/Code/OpenACC/OpenACCProject/OpenACCProject$ ./main_acc_ubuntu.exe
histgram[0] = 123
histgram[1] = 79
histgram[2] = 90
histgram[3] = 109
histgram[4] = 102
histgram[5] = 96
histgram[6] = 117
histgram[7] = 104
histgram[8] = 110
histgram[9] = 94

● 添加一个中间变量用来存储 data[i]，在 Windos 下也正确了

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

int main()
{
    const int length = 1024, basket = 10;
    int data[length], histgram[basket];

    srand(97);
    for (int i = 0; i < basket; histgram[i++] = 0);
    for (int i = 0; i < length; data[i++] = rand() % 10);

    int temp;               // 新定义一个变量
#pragma acc parallel loop
    for (int i = 0; i < length; i++)
    {
        temp = data[i];     // 原子操作之前先单独计算下标
#pragma acc atomic update
        histgram[temp] += 1;// 使用 temp 作为下标
    }

    for (int i = 0; i < basket; i++)
        printf("histgram[%d] = %d
", i, histgram[i]);

    getchar();
    return 0;
}