如何使用 CUDA 并行化嵌套 for 循环以对 2D 数组执行计算

Question

我正在进行一些研究，并且是使用 CUDA 的初学者。我使用的语言是 C 和 C++，与 Nvidia 的 CUDA 兼容的基本语言。在过去的一周里，我一直试图通过将 CUDA 与我的 C++ 代码集成来获得任何形式的加速。

据我所知，就内存分配和释放而言，我做的基本都是正确的。但是当谈到实际加速计算时，我目前从非 CUDA 实现中收到不同的结果。

此外，CUDA 实现也比普通的非 cuda 版本慢。

以下是我从中调用内核函数的函数。本质上，我将原本在这个函数中的计算移到了内核函数中，以便并行化它。 //计算输入之间的距离 void computeInput(int vectorNumber, double *dist, double **weight){

double *d_dist, **d_weight;


//cout << "Dist[0] Before: " << dist[0] << endl;

cudaMalloc(&d_dist, maxClusters * sizeof(double));
cudaMalloc(&d_weight, maxClusters * vector_length * sizeof(double));

//  cout << "Memory Allocated" << endl;

//copy variables from host machine running on CPU to Kernel running on GPU
cudaMemcpy(d_dist, dist, maxClusters * sizeof(double), cudaMemcpyHostToDevice);
cudaMemcpy(d_weight, weight, maxClusters * vector_length * sizeof(double), cudaMemcpyHostToDevice);

//  cout << "Variables copied to GPU Device." << endl;

//kernel currently being run with 1 blocks with 4 threads for each block.
//right now only a single loop is parallelized, I need to parallelize each loop individually or 2d arrays individually.
dim3 blocks(8,8);
dim3 grid(1, 1);
threadedInput<<<grid,blocks>>>(vectorNumber, d_dist, d_weight);

//  cout << "Kernel Run." << endl;  

//Waits for the GPU to finish computations
cudaDeviceSynchronize();

//cout << "Weight[0][0] : " << weight[0][0];

//copy back varaible from kernelspace on GPU to host on CPU into variable weight
cudaMemcpy(weight, d_weight, maxClusters * vector_length * sizeof(double), cudaMemcpyDeviceToHost);
cudaMemcpy(dist, d_dist, maxClusters * sizeof(double), cudaMemcpyDeviceToHost);
//  cout << "GPU Memory Copied back to Host" << endl;

cout << "Dist[0] After: " << dist[0] << endl;

cudaFree(d_dist);
cudaFree(d_weight);

//cout << " Cuda Memory Freed" << endl;
}

以下是内核函数。它正在使用节点上的权重计算距离。

我想要它做的是在单独的线程上执行循环的每次迭代。

我担心它会弄乱顺序并执行错误的计算。我已经在 Stack Overflow 和其他地方搜索了有关嵌套 for 循环并行化的帮助，但它们都没有说明我做错了什么。有什么建议吗？

__global__ void threadedInput(int vecNum, double *dist, double **weight)
{
int tests[vectors][vector_length] = {{0, 1, 1, 0},
                                     {1, 0, 0, 1},
                                     {0, 1, 0, 1},
                                     {1, 0, 1, 0}};
dist[0] = 0.0;
dist[1] = 0.0;
int indexX,indexY, incrX, incrY;
indexX = blockIdx.x * blockDim.x + threadIdx.x;
indexY = blockIdx.y * blockDim.y + threadIdx.y;
incrX = blockDim.x * gridDim.x; 
incrY = blockDim.y * gridDim.y; 

for(int i = indexY; i <= (maxClusters - 1); i+=incrY)
{
    for(int j = indexX; j <= (vectors - 1); j+= incrX)
    {       
        dist[i] += pow((weight[i][j] - tests[vecNum][j]), 2);
    }// end inner for
}// end outer for

}// end CUDA-kernel

我当前的输出：

Clusters for training input:

Vector (1, 0, 1, 0, ) Place in Bin 0

Vector (1, 1, 1, 0, ) Place in Bin 0

Vector (0, 1, 1, 1, ) Place in Bin 0

Vector (1, 1, 0, 0, ) Place in Bin 0

Weights for Node 0 connections:
0.74753098, 0.75753881, 0.74233157, 0.25246902, 

Weights for Node 1 connections:
0.00000000, 0.00000000, 0.00000000, 0.00000000, 

Categorized test input:

Vector (0, 1, 1, 0, ) Place in Bin 0

Vector (1, 0, 0, 1, ) Place in Bin 0

Vector (0, 1, 0, 1, ) Place in Bin 0

Vector (1, 0, 1, 0, ) Place in Bin 0
Time Ran: 0.96623900

预期输出（除了预期的时间应该至少快 50%）

Clusters for training input:

Vector (1, 0, 1, 0, ) Place in Bin 0

Vector (1, 1, 1, 0, ) Place in Bin 1

Vector (0, 1, 1, 1, ) Place in Bin 0

Vector (1, 1, 0, 0, ) Place in Bin 1

Weights for Node 0 connections:
0.74620975, 0.75889148, 0.74351981, 0.25379025, 

Weights for Node 1 connections:
0.75368531, 0.75637331, 0.74105526, 0.24631469, 

Categorized test input:

Vector (0, 1, 1, 0, ) Place in Bin 0

Vector (1, 0, 0, 1, ) Place in Bin 1

Vector (0, 1, 0, 1, ) Place in Bin 0

Vector (1, 0, 1, 0, ) Place in Bin 1
Time Ran: 0.00033100

Answer 1

你应该阅读一些教程，开头是：https://devblogs.nvidia.com/easy-introduction-cuda-c-and-c/

基本上每个线程都执行内核代码，所以里面应该没有循环。

我在引用：

设备代码

我们现在转到内核代码。

__global__
void saxpy(int n, float a, float *x, float *y)
{
  int i = blockIdx.x*blockDim.x + threadIdx.x;
  if (i < n) y[i] = a*x[i] + y[i];
}

在 CUDA 中，我们使用 global 声明 >specifier 定义内核，例如 saxpy。在设备代码中定义的变量不需要指定为设备变量，因为它们被假定驻留在设备上。在这种情况下，n、a 和 i 变量将由每个线程存储在 > 寄存器中，并且指针 x 和 y 必须是指向设备内存 > 地址空间的指针。这确实是真的，因为当我们从主机代码启动它时，我们将 d_x 和 d_y 传递给 >kernel。然而，前两个参数 n > 和 a 并未在主机代码中显式传输到设备。 >由于函数参数在 C/C++ 中默认按值传递，因此 >CUDA 运行时可以自动处理将这些值传输到 >device。 CUDA Runtime API 的这一特性使得在 >GPU 上启动内核变得非常自然和容易——它几乎与调用 C 函数一样。

在我们的 saxpy 内核中只有两行。如前所述，>kernel 是由多个线程并行执行的。如果我们希望每个线程>处理结果数组的一个元素，那么我们需要一种>区分和识别每个线程的方法。 CUDA 定义了变量 >blockDim、blockIdx 和 threadIdx。这些预定义变量的类型为 >dim3，类似于主机代码中的执行配置参数。 >predefined 变量 blockDim 包含每个线程块的尺寸 > 如在内核 >launch 的第二个执行配置参数中指定的。预定义变量threadIdx 和blockIdx 分别包含其线程块内的线程和网格内的线程块的索引>。表达式：

int i = blockDim.x * blockIdx.x + threadIdx.x

生成用于访问数组元素的全局索引。我们>在此示例中没有使用它，但还有 gridDim 包含>在第一个执行配置中指定的网格尺寸>启动参数。

在使用此索引访问数组元素之前，会根据元素的数量 n 检查其值，以确保没有超出范围的内存访问。如果数组中的 > 元素数不能被线程块大小整除，并且 > 因此内核启动的线程数大于 > 数组大小，则需要进行此检查。内核的第二行执行 >SAXPY 的逐元素工作，除了边界检查之外，它与 SAXPY 的主机实现的内部 >loop 相同。

if (i < n) y[i] = a*x[i] + y[i];