【发布时间】：2016-02-19 04:30:33
【问题描述】：
我对理解SM中的CUDA线程处理有些疑惑。从我一直在阅读的内容中推断出以下命题： 我的 GPU 是：GTX650Ti。
- 块中的线程计数必须始终是 Warp 大小的倍数。因此，每个 SM 可以处理 32 个线程的块 (warpSize)。
- 我的 SM 可以同时计算的最大线程数为 2048 (maxThreadsPerMultiProcessor)。
- 由于每个 SM 可以同时计算 2048 个线程，warpSize 为 32，因此可以同时计算 64 个块。
- 由于我的GPU有4个SM，可以同时执行64X4=256个线程块。
- 因此，内核启动可能有以下启动参数：>>，每次内核启动会调用8192个线程。
是吗？
因此，如果我的内核中有一个包含 10M 元素的向量要处理，这意味着我必须将其分割为 1221 个作业（内核启动），每个作业包含 8192 个元素？
之所以出现这个问题，是因为我正在比较顺序程序和我的 CUDA 程序之间的时间性能。但我只能看到 CPU 超过了 GPU。我还尝试了最大启动参数，例如 >>。结果非常相似。
那么，我在做什么或配置错误？
这是我正在使用的代码：
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <math.h>
#include <time.h>
#include "C:\cdev.h"
#include <thrust/device_vector.h>
using namespace thrust;
using namespace std;
#define N (1024 * 16384)
cdev devices;
__global__ void eucliDist(double *c, double *a, double *b)
{
int i = blockDim.x * blockIdx.x + threadIdx.x;
if (i < N)
c[i] = sqrt(pow(a[i], 2) + pow(b[i], 2));
}
int main()
{
clock_t start, end;
double elapsed;
static double A[N];
static double B[N];
for (int i = 0; i < N; i++)
{
A[i] = double(i);
B[i] = double(i);
}
static double C[N];
// Sequential execution of F(x,y) = sqrt((x^2 + y^2))
start = clock();
for (int i = 0; i < N; i++)
C[i] = sqrt(pow(A[i], 2) + pow(B[i], 2));
end = clock();
elapsed = double(end - start) / CLOCKS_PER_SEC;
cout << "Elapsed time for sequential processing is: " << elapsed << " seconds." << endl;
// CUDA Initialization
unsigned int threadNum;
unsigned int blockNum;
cudaError_t cudaStatus;
threadNum = devices.ID[0].maxThreadsPerBlock;
blockNum = ceil(double(N) / double(threadNum));
// Parallel execution with Thrust of F(x,y) = sqrt((x^2 + y^2))
vector<double> vectorA(N);
vector<double> vectorB(N);
for (int i = 0; i < N; i++)
{
vectorA[i] = double(i);
vectorB[i] = double(i);
}
vector<double> vectorC(N);
start = clock();
device_vector<double> thrustA(N);
cudaStatus = cudaGetLastError();
if (cudaStatus != cudaSuccess)
{
cerr << "Device vector allocation failed: " << cudaGetErrorString(cudaStatus) << " (thrustA)" << endl;
cin.get();
return 1;
}
device_vector<double> thrustB(N);
cudaStatus = cudaGetLastError();
if (cudaStatus != cudaSuccess)
{
cerr << "Device vector allocation failed: " << cudaGetErrorString(cudaStatus) << " (thrustB)" << endl;
cin.get();
return 1;
}
device_vector<double> thrustC(N);
cudaStatus = cudaGetLastError();
if (cudaStatus != cudaSuccess)
{
cerr << "Device vector allocation failed: " << cudaGetErrorString(cudaStatus) << " (thrustC)" << endl;
cin.get();
return 1;
}
thrustA = vectorA;
cudaStatus = cudaGetLastError();
if (cudaStatus != cudaSuccess)
{
cerr << "Host to device copy failed (Thrust): " << cudaGetErrorString(cudaStatus) << " (vectorA -> thrustA)" << endl;
cin.get();
return 1;
}
thrustB = vectorB;
cudaStatus = cudaGetLastError();
if (cudaStatus != cudaSuccess)
{
cerr << "Host to device copy failed (Thrust): " << cudaGetErrorString(cudaStatus) << " (vectorB -> thrustB)" << endl;
cin.get();
return 1;
}
eucliDist <<<blockNum, threadNum>>>(raw_pointer_cast(thrustC.data()), raw_pointer_cast(thrustA.data()), raw_pointer_cast(thrustB.data()));
cudaStatus = cudaGetLastError();
if (cudaStatus != cudaSuccess)
{
cerr << "Kernel launch failed (Thrust): " << cudaGetErrorString(cudaStatus) << " (euclidDist)" << endl;
cin.get();
return 1;
}
thrust::copy(thrustC.begin(), thrustC.end(), vectorC.begin());
cudaStatus = cudaGetLastError();
if (cudaStatus != cudaSuccess)
{
cerr << "Device to host copy failed: " << cudaGetErrorString(cudaStatus) << " (thrustC -> vectorC)" << endl;
cin.get();
return 1;
}
end = clock();
elapsed = double(end - start) / CLOCKS_PER_SEC;
cout << "Elapsed time parallel processing is (Thrust): " << elapsed << " seconds." << endl;
cin.get();
return 0;
}
建议将不胜感激。
【问题讨论】：
-
在分析器中运行你的程序，看看时间花在了哪里。不要挂在硬件细节上，只需启动内核一次，使用尽可能多的块来覆盖整个 10M 元素的网格。从教科书“添加向量”示例开始学习该技术。而且 GPU 并不总是比 CPU 快；这取决于任务。
-
我会说，在第 1 点的第一句话之后，您列表中的所有内容都不正确...
-
我可以证明这一点。请参阅 CUDA 编程：Shane Cook 的 GPU 并行计算开发人员指南，第 83 页，GRIDS 部分，第 2 段，第 1 行：“块中的线程数应始终是 Warp 大小的倍数，目前定义为 32。”当我说必须时，我可能夸大了，所以这只是建议不要使用条件语句来处理向量大小的元素。
-
您在很多方面感到困惑，并有效地要求在这个问题中提供有关许多 CUDA 主题的广泛教程。正如@talonmies 所说，您几乎所有的猜想都是错误的。从哪儿开始？也许你应该研究一个简单的代码，比如 vectorAdd 并观察它可以在单个内核启动中处理一个大向量。您绝对不需要 1221 内核启动来处理 10M 元素向量。
-
@Vitrion：但为什么它应该自动更快呢？您完全没有提供有关您的代码的任何信息，因此无法说出您的代码可能存在什么性能问题或您可能会采取什么措施来修复它们。