1

我需要关于如何继续和利用 CUDA 设备的计算能力来对函数进行数值积分的建议。关于我的设备的一些信息如下（不相关）

硬件

 Geforce GTX470; Compute Capability 2.0

问题描述

我有一个像

g(x) = x * f(x, a, b, c)

我需要按照给定的方程进行积分

现在我已经编写了一个积分函数，它只需要 g(x)，将区间分成N个子区间，计算各个子区间的结果，然后在 CPU 上进行汇总。为了完成目的，我在下面提供了一个代码示例。

__device__ float function(float x, float a, float b, float c) {
   // do some complex calculation
   return result;
}
__global__ void kernel(float *d_arr, float a, float b, float c, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    float x = (float)idx / (float)N;

    if (idx < N)  {
       d_arr[idx] = x * function(x, a, b, c);
    }
}

上面的代码仅用于演示目的，我实际上使用 Romberg 方法来集成我的g(x)，但想法是相同的。我真正的问题是因为我不只有一组值（a，b，c），我有这组的多个值。

我在设备内存中有一个二维数组，正好是 (3, 1024) 3 行，1024 列。每列代表一个需要执行集成功能的集合。

当我必须决定是否执行一个线程块（例如 1024）时，问题就来了，记住一个线程相当于一个集成函数。在这种情况下，我上面写的函数是没有用的。因为我想对所有的值集进行并行积分，所以我必须编写一个积分函数，它可以按顺序进行积分。举个例子：

__global__ void kernel(float *d_arr, float a, float b, float c, int N) {
    
   int idx = blockIdx.x * blockDim.x + threadIdx.x;
   float sum = 0;
   for (int i = 0; i < N; i++) {
      float x = (float)i / (float) N;
      sum += x * function(x, a, b, c);
   } 
    d_arr[idx] = sum;
}

所以你明白我的意思了吗？选项 A 似乎更好，但我不能使用它，因为我不知道如何进行多个积分，然后将每个积分分配给 N 个线程。

你会怎么做？你能建议我，我怎样才能实现多个积分，而每个积分都可以分布到 N 个线程？有没有更好的方法来做到这一点。

期待您的建议。

4

1 回答 1

1

如果我正确理解您的问题，您想对多组（1024）输入（a，b，c）进行数值积分，并且对于每个积分，您需要 N 个子区间。我们称输入集合的数量为 M。

如果 N 足够大（假设 > 10000），您粘贴的第一个内核样本可能就足够好（针对不同的输入集调用它 M 次）。它是否利用所有可用的设备吞吐量取决于您的功能有多复杂。

我不明白你对 d_arr[] 数组到底做了什么？通常对于数值积分，您需要对其求和。正确的？你是在总结 CPU 上的结果吗？如果您发现 atomicAdd 不够快，请考虑使用 atomicAdd（尤其是如果您要在计算 cap 3.0 及更高版本的 gpus 上运行内核）或并行扫描。

如果 N 很小，最好在单个内核中启动 N*M 个线程。

在您的 M=1024 的情况下，您可以让每个块处理一组输入（即设置 blockSize = 1024），并将 (a,b,c) 输入作为数组传递给内核 - 如下所示：

__global__ void kernel(float *d_arr, float *a_array, float *b_array, float *c_array, int totalThreads, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    float x = (float) blockIdx.x / (float) N;
    float a = a_array[threadIdx.x];
    float b = b_array[threadIdx.x];
    float c = c_array[threadIdx.x];

    if (idx < totalThreads)  {
       // what happen to this array?
       d_arr[idx] = x * function(x, a, b, c);
    }
}

同样，您稍后需要从 d_arr 从适当的位置提取元素并将它们相加（对于每个积分）。

如果您的函数不是很复杂并且上述内核受内存限制，您可以尝试相反的方式，即让每个线程块处理每个子间隔 - 不同的线程块处理不同的输入集。内核看起来像这样：

（此示例假设 N <= 1024，但即使不是，也可以分解内核以利用这种方法）

__global__ void kernel(float *d_arr, float *a_array, float *b_array, float *c_array, int totalThreads) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;

    float x = (float)threadIdx.x / (float) blockDim.x;  // N = blockDim.x

    float a = a_array[blockIdx.x];  // every thread in block accesses same memory location
    float b = b_array[blockIdx.x];
    float c = c_array[blockIdx.x];

    // d_arr has 'M' elements containing the integral for each input set.
    if (idx < totalThreads)  
    {
       atomicAdd(&d_arr[blockIdx.x], x * function(x, a, b, c));
    }
}

在上述内核中，a_array、b_array 和 c_array 分配在常量内存中。这会更快，因为块中的每个线程都将访问相同的位置。例如，我还用 atomicAdd 替换了您的 d_arr 写入。

于 2013-06-16T19:40:58.287 回答