如何区分单个 GPU 中不同主机 CPU 线程的 GPU 线程
how to differentiate GPU threads in a single GPU for different host CPU thread
当多个 CPU 线程将作业分派给单个 GPU 时，区分 GPU 线程的最佳方法是什么，以便多个 CPU 线程不会简单地相互重复
以下代码逐个元素计算两个大数组的总和。正确的结果是：3.0。当使用 1 CPU 时，代码做正确的事情。然后运行8CPUs，输出变为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
这个非常简单的案例的解决方案是将数组分成几部分，每个线程一个。为简单起见，这样我就不必处理一堆恼人的极端情况问题，让我们假设您的数组大小 (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 线程并没有什么特别的好处。我假设这里要问的是设计模式以启用其他 activity。多个 CPU 线程可能是安排其他工作的便捷方式。例如，我可能有一个系统正在处理来自 4 个摄像头的数据。将我的相机处理组织为 4 个独立线程可能很方便，每个线程一个。该系统可能只有 1 个 GPU，4 个线程中的每一个都可能希望向该 GPU 发出独立的工作，这当然是合理的。举个例子，这种设计模式可以很容易地适应那个用例。甚至可能是 4 个摄像头 CPU 线程需要将一些数据组合到 GPU 上的单个数组中，在这种情况下可以使用这种模式。
When multiple CPU thread dispatch jobs to a single GPU, what's the best way to differentiate GPU threads so that the multiple CPU thread does not simply repeat each other
让我比你的具体例子更笼统地回答这个问题：
- 使用多线程在 GPU 上排队工作没有内在的好处。如果您让每个线程都在一个 CUDA 队列上等待，那么它可能有意义，但这不一定是正确的做法。
- 除非您明确安排内存传输，否则无法保证将您安排的工作分成小块的固有好处。您可以只安排一个内核来添加整个阵列。请记住——一个内核在 GPU 端由数千或数百万个 'threads' 组成； CPU 线程对 GPU 并行性根本没有帮助。
- 当不同的线程开始意识到它们彼此独立存在时，让它们安排工作更有意义。
- 将内核的输出写在与其输入不同的地方通常是个好主意。它在计算过程中需要更多内存，但它可以防止您描述的那种问题 - 相同值的重叠更改，必须仔细考虑哪个调度内核首先执行等。因此，例如，您可以实现：
__global__ void add_2(int n, float* result, const float *x, const float *y)
{
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) {
z[i] = x[i] + y[i];
}
}
如果你不能这样做，那么你需要仔细划分输入输出数组来安排工作，正如@RobertCrovella 的回答中所建议的那样。
使用 __restrict__ 关键字（即使它不是标准的 C++）来指示参数指向的区域不重叠。这加快了速度。参见：
当多个 CPU 线程将作业分派给单个 GPU 时，区分 GPU 线程的最佳方法是什么，以便多个 CPU 线程不会简单地相互重复
以下代码逐个元素计算两个大数组的总和。正确的结果是：3.0。当使用 1 CPU 时，代码做正确的事情。然后运行8CPUs，输出变为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
这个非常简单的案例的解决方案是将数组分成几部分，每个线程一个。为简单起见，这样我就不必处理一堆恼人的极端情况问题，让我们假设您的数组大小 (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 线程并没有什么特别的好处。我假设这里要问的是设计模式以启用其他 activity。多个 CPU 线程可能是安排其他工作的便捷方式。例如，我可能有一个系统正在处理来自 4 个摄像头的数据。将我的相机处理组织为 4 个独立线程可能很方便，每个线程一个。该系统可能只有 1 个 GPU，4 个线程中的每一个都可能希望向该 GPU 发出独立的工作，这当然是合理的。举个例子，这种设计模式可以很容易地适应那个用例。甚至可能是 4 个摄像头 CPU 线程需要将一些数据组合到 GPU 上的单个数组中，在这种情况下可以使用这种模式。
When multiple CPU thread dispatch jobs to a single GPU, what's the best way to differentiate GPU threads so that the multiple CPU thread does not simply repeat each other
让我比你的具体例子更笼统地回答这个问题：
- 使用多线程在 GPU 上排队工作没有内在的好处。如果您让每个线程都在一个 CUDA 队列上等待，那么它可能有意义，但这不一定是正确的做法。
- 除非您明确安排内存传输，否则无法保证将您安排的工作分成小块的固有好处。您可以只安排一个内核来添加整个阵列。请记住——一个内核在 GPU 端由数千或数百万个 'threads' 组成； CPU 线程对 GPU 并行性根本没有帮助。
- 当不同的线程开始意识到它们彼此独立存在时，让它们安排工作更有意义。
- 将内核的输出写在与其输入不同的地方通常是个好主意。它在计算过程中需要更多内存，但它可以防止您描述的那种问题 - 相同值的重叠更改，必须仔细考虑哪个调度内核首先执行等。因此，例如，您可以实现：
如果你不能这样做，那么你需要仔细划分输入输出数组来安排工作，正如@RobertCrovella 的回答中所建议的那样。__global__ void add_2(int n, float* result, const float *x, const float *y) { int i = blockIdx.x * blockDim.x + threadIdx.x; if (i < n) { z[i] = x[i] + y[i]; } } 使用
__restrict__关键字（即使它不是标准的 C++）来指示参数指向的区域不重叠。这加快了速度。参见：