c++ - 如何在 CUDA 中实现基本的 C++ 二维数组循环

Question

我是 C++ 编码的新手，目前正在尝试使用 CUDA 进行一些 GPU 计算。

基本上我有一个矩阵A(N乘N)，还有几个 vector b和x0。 b 和 x0 也有 N 个元素。

这是我要实现的一段代码:

for (unsigned i=1;i<=N;i++){
    T sum = 0;
    for (unsigned j=1;j<=N;j++){
        sum += A[j][i]*x0[j];
    }
    v[i] = b[i] - sum;
}

其中 T 是一个模板变量(据我所知可以分配一个 double 值)。

是否有可能将整个事情并行化，如果可以，我会怎么做？我还可以使用一些关于如何将此类问题的线程分解为 block 以及如何将 2D 从主机移动到设备并返回......

如果需要任何其他信息，请告诉我。

编辑 1: 在研究了 CUBLAS 并没有走得太远之后，我决定展平我的矩阵并自己编写代码。我的第一个发现是我的 cuda 内核不喜欢使用 double 类型变量/数组 [有人可以证实这一点吗？]。

在将所有内容都转换为 float 之后，我编写的 cuda 内核看起来像这样:

__global__ void cudaMatTimesVect(float *p, float  *x0, float *v, float *sum, float *toSum, float *b, int N){

int idx = blockIdx.x * blockDim.x + threadIdx.x; // thread index

if (idx < N*N){
    toSum[idx] = p[idx] * x0[blockIdx.x];
}

__syncthreads();
if( idx-(blockIdx.x * blockDim.x) == 0){
    for(int i=0; i<blockDim.x; i++){
        sum[blockIdx.x] += toSum[idx+i];
    }

v[blockIdx.x] = b[blockIdx.x] - sum[blockIdx.x];
}

我不确定 syncthreads() 命令是否会在尝试执行求和循环之前等待所有线程相乘。

这里是关于仅在 GPU 上初始化的 sum 和 toSum 数组的 CPU 代码片段:

float *d_sum;
float *d_toSum;
cudaError_t  cudaStatus;
...
// allocate toSum memory
cudaStatus = cudaMalloc(&d_toSum, N*N*sizeof(float));
if (cudaStatus != cudaSuccess){
    std::cout << "couldnt allocate device memory for d_toSum!" << std::endl;
    cudaFree(d_toSum);
}
// allocate sum mem on device
cudaStatus = cudaMalloc(&d_sum, N*sizeof(float));
if (cudaStatus != cudaSuccess){
    std::cout << "couldnt allocate device memory for d_sum" << std::endl;
    cudaFree(d_sum);
}

... 
...
// call the kernel
cudaMatTimesVect<<<N,N>>>(d_p, d_x0, d_v, d_sum, d_toSum, d_b, N);
...


cudaFree(d_toSum);
cudaFree(d_sum);

这是求和的最有效方法吗？

编辑 2: 我现在更改了代码以使用不同的 block 索引来运行行计算。 上面的内核编译运行，但是v中的数组元素似乎越来越小而不是重启...

我仍然有兴趣了解为什么我不能使用 double ，以及如果我想使用 定义我的主机数组，我的代码需要如何更改。

谢谢，

亚美尼亚人

Answer 1

你可以在cublas中解决这个问题:

使用 cublasSetVector 将数据复制到 GPU或 cublasSetMatrix

使用相应的Get将结果复制回来functions .

矩阵 vector 乘法由 gemv 处理. vector - vector 减法用 axpy 处理。 .

cuda samples 中提供了一个有效的 cublas 示例.

根据附加评论: 对于这个问题，没有理由将数据分割成一维 block 。我推荐cublas。但如果您想查看其他代码示例，请查看 vector add example和 matrix multiply example .

对于主机上的双下标矩阵，您应该将其展平，以便您可以使用单个 (*) 指针和索引来引用数据。无论您使用的是 cublas 还是编写自己的代码，都是如此。

编辑:响应问题中的更新。 您发布的乘法代码对我来说看起来不像矩阵 vector 乘法，除非您将 vector 的长度复制 N 次以使其与矩阵的长度 (NxN) 匹配。那么它似乎是正确的。

求和代码看起来不正确，而且，由于它不以任何方式依赖于 idx，所有线程 都在做完全相同的事情.所以那里没有并行的好处，我们通常不会以这种方式编写 GPU 代码。

你的 vector 减法代码似乎是大致正确的，除了你似乎在矩阵的整个长度 (NxN) 上做 vector 减法，而矩阵 vector 乘法的结果应该只产生一个 vector 长度N。

如果此代码生成的结果与相同数据集的序列代码相匹配，我会感到很惊讶。您是否检查过它是否为非平凡的数据集产生了正确的结果？ (不要使用每个数字都相同的数据集。)