c++ - 如何区分单个 GPU 中不同主机 CPU 线程的 GPU 线程

Question

当多个 CPU 线程向单个 GPU 分配作业时，区分 GPU 线程的最佳方式是什么，以便多个 CPU 线程不会简单地相互重复

以下代码逐个元素地计算两个大数组的总和。正确的结果是:3.0。当使用 1 个 CPU 时，代码会做正确的事情。然后用 8 个 CPU 运行，输出变为 10，因为内核重复计算 8 次。我正在寻找一种方法，使每个 CPU 计算不相互重复的总和的 1/8。

#include <iostream>
#include <math.h>
#include <thread>
#include <vector>

#include <cuda.h>


using namespace std;

const unsigned NUM_THREADS = std::thread::hardware_concurrency();  

// Kernel function to add the elements of two arrays
__global__
void add_2(int n, float *x, float *y)
{
    int i = blockIdx.x*blockDim.x + threadIdx.x;
    if(i < n) {
        y[i] = x[i] + y[i];
    }
}

//
void thread_func(int N, float *x, float *y, int idx_thread)
{   
    cudaSetDevice(0); 

    int blockSize;   
    int minGridSize; 
    int gridSize;    

    cudaOccupancyMaxPotentialBlockSize( &minGridSize, &blockSize, add_2, 0, N);
    // Round up according to array size
    gridSize = (N + blockSize - 1) / blockSize;
    //gridSize /= NUM_THREADS +1;  

    cout<<"blockSize: "<<blockSize<<" minGridSize: "<<minGridSize<<" gridSize: "<<gridSize<<endl;

    // Run kernel on 1M elements on the GPU
    add_2<<<gridSize, blockSize>>>(N, x, y);


    // Wait for GPU to finish before accessing on host
    cudaDeviceSynchronize();
}


//
int main()
{

    int N = 1<<20;
    float *x, *y;
    // Allocate Unified Memory – accessible from CPU or GPU
    cudaMallocManaged(&x, N*sizeof(float));
    cudaMallocManaged(&y, N*sizeof(float));

    // initialize x and y arrays on the host
    for (int i = 0; i < N; i++) {
        x[i] = 1.0f;
        y[i] = 2.0f;
    }

    //.. begin multithreading ..
    vector<std::thread> t;
    for(int i = 0; i<NUM_THREADS; i++)
        t.push_back(thread(thread_func, N, x, y, i));

    for(int i = 0; i<NUM_THREADS; i++)
        t[i].join();

    // Check for errors (all values should be 3.0f)
    float maxError = 0.0f;
    for (int i = 0; i < N; i++) {
        if(!(i%10000))
            std::cout<<i<<" "<<y[i]<<std::endl;
        maxError = fmax(maxError, fabs(y[i]-3.0f));
    }
    std::cout << "Max error: " << maxError << std::endl;

    // Free memory
    cudaFree(x);
    cudaFree(y);

    return 0;
}

输出:

blockSize: 1024 minGridSize: 16 gridSize: 1024

..........

blockSize: 1024 minGridSize: 16 gridSize: 1024

0 10

10000 10

20000 10

...

1020000 10

1030000 10

1040000 10

Max error: 7

Answer 1

这个非常简单的案例的解决方案是将数组分成几部分，每个线程一个。为简单起见，这样我就不必处理一堆恼人的极端情况问题，让我们假设您的数组大小 (N) 是可被 NUM_THREADS 整除的整数。当然，不一定非要这样，但划分它的算法并没有太大不同，但你必须处理每个段边界的舍入，我宁愿避免这种情况。

这是一个基于上述假设的示例。每个线程决定它负责数组的哪一部分(基于它的线程数和总长度)并且只在该部分上工作。

$ cat t1460.cu
#include <iostream>
#include <math.h>
#include <thread>
#include <vector>

#include <cuda.h>


using namespace std;

const unsigned NUM_THREADS = 8;

// Kernel function to add the elements of two arrays
__global__
void add_2(int n, float *x, float *y)
{
    int i = blockIdx.x*blockDim.x + threadIdx.x;
    if(i < n) {
        y[i] = x[i] + y[i];
    }
}

//
void thread_func(int N, float *x, float *y, int idx_thread)
{
    cudaSetDevice(0);

    int blockSize = 512;
    int worksize = N/NUM_THREADS; // assumes whole-number divisibility
    int gridSize = (worksize+blockSize-1)/blockSize;
    cout<<"blockSize: "<<blockSize<<" gridSize: "<<gridSize<<endl;

    // Run kernel on 1M elements on the GPU
    add_2<<<gridSize, blockSize>>>(worksize, x+(idx_thread*worksize), y+(idx_thread*worksize));


    // Wait for GPU to finish before accessing on host
    cudaDeviceSynchronize();
}


//
int main()
{

    int N = 1<<20;
    float *x, *y;
    // Allocate Unified Memory – accessible from CPU or GPU
    cudaMallocManaged(&x, N*sizeof(float));
    cudaMallocManaged(&y, N*sizeof(float));

    // initialize x and y arrays on the host
    for (int i = 0; i < N; i++) {
        x[i] = 1.0f;
        y[i] = 2.0f;
    }

    //.. begin multithreading ..
    vector<std::thread> t;
    for(int i = 0; i<NUM_THREADS; i++)
        t.push_back(thread(thread_func, N, x, y, i));

    for(int i = 0; i<NUM_THREADS; i++)
        t[i].join();

    // Check for errors (all values should be 3.0f)
    float maxError = 0.0f;
    for (int i = 0; i < N; i++) {
        if(!(i%10000))
            std::cout<<i<<" "<<y[i]<<std::endl;
        maxError = fmaxf(maxError, fabs(y[i]-3.0f));
    }
    std::cout << "Max error: " << maxError << std::endl;

    // Free memory
    cudaFree(x);
    cudaFree(y);

    return 0;
}
$ nvcc t1460.cu -o t1460 -std=c++11
$ cuda-memcheck ./t1460
========= CUDA-MEMCHECK
blockSize: blockSize: 512 gridSize: 256512blockSize:  gridSize:
blockSize: blockSize: 512blockSize:  gridSize: 256512
 gridSize: 256
blockSize: 512 gridSize: 256
blockSize: 512 gridSize: 256
512 gridSize: 256
256
512 gridSize: 256
0 3
10000 3
20000 3
30000 3
40000 3
50000 3
60000 3
70000 3
80000 3
90000 3
100000 3
110000 3
120000 3
130000 3
140000 3
150000 3
160000 3
170000 3
180000 3
190000 3
200000 3
210000 3
220000 3
230000 3
240000 3
250000 3
260000 3
270000 3
280000 3
290000 3
300000 3
310000 3
320000 3
330000 3
340000 3
350000 3
360000 3
370000 3
380000 3
390000 3
400000 3
410000 3
420000 3
430000 3
440000 3
450000 3
460000 3
470000 3
480000 3
490000 3
500000 3
510000 3
520000 3
530000 3
540000 3
550000 3
560000 3
570000 3
580000 3
590000 3
600000 3
610000 3
620000 3
630000 3
640000 3
650000 3
660000 3
670000 3
680000 3
690000 3
700000 3
710000 3
720000 3
730000 3
740000 3
750000 3
760000 3
770000 3
780000 3
790000 3
800000 3
810000 3
820000 3
830000 3
840000 3
850000 3
860000 3
870000 3
880000 3
890000 3
900000 3
910000 3
920000 3
930000 3
940000 3
950000 3
960000 3
970000 3
980000 3
990000 3
1000000 3
1010000 3
1020000 3
1030000 3
1040000 3
Max error: 0
========= ERROR SUMMARY: 0 errors
$

当然，对于这个简单的示例，使用 4 个 CPU 线程并没有特别的好处。我假设这里要求的是一种设计模式来启用其他事件。多个 CPU 线程可能是安排其他工作的便捷方式。例如，我可能有一个系统正在处理来自 4 个摄像头的数据。将我的相机处理组织为 4 个独立线程可能很方便，每个线程一个。该系统可能只有 1 个 GPU，4 个线程中的每一个都可能希望向该 GPU 发出独立的工作，这当然是合理的。举个例子，这种设计模式可以很容易地适应那个用例。甚至可能是 4 个相机 CPU 线程需要将一些数据组合到 GPU 上的单个数组中，在这种情况下可以使用这种模式。