-1

我试图通过将来自命令行参数的值传递给内核来使 Cuda 应用程序更具动态性。

该应用程序调用多个内核，并最大化块和网格大小。当我尝试运行应用程序时，我得到以下结果：

  • 硬编码值：0.96 秒
  • 在内核初始化时传递一个值：3.48 秒
  • 声明 a __device__ int，并将其设置为值：3.48 秒

一旦在执行时输入该值，它将在程序的其余部分保持不变。

这两个 3.48 秒时间来自对变量本身的访问。如果我用硬编码的整数替换变量，运行时间会大大缩短。这个值被非常频繁地访问，我想知道是否有办法保持与硬编码值相似的速度，但降低访问变量的成本。是否可以通过使用变量来加快速度？

慢 3.6 倍重要吗？有点。这只是一小部分更大的东西。

任何帮助将不胜感激。

*运行 2.0 硬件。

编辑：这是我遇到的差异的一个例子：

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
#include <ctime>

using namespace std;

clock_t start;

__device__ int x;

__global__ void setNum(int i)
{
        x = i;
        return;
}

__device__ void d_swap(int * a, int * b)
{
        int temp = *a;
        *a = *b;
        *b = temp;
}

__device__ void other(int n, int * vec)
{
        int i;
        for(i = 0; i < n; ++i) vec[i] = i;
        for (int j = 0; j < 5; j++)
                for(i = 1; i < n-1; ++i)
                        d_swap(&vec[i], &vec[i-1]);
}

__global__ void Pressure(int i)
{
        int a[12];
        other(x, a);
        //other(12,a);
}

int main(int argc, char * argv[])
{
        if (argc != 2)
        {
                fprintf(stderr,"Invalid number of arguments.\n");
                exit(1);
        }
        int num = atoi(argv[1]);
        cudaSetDevice(1);

        cudaMemset(&x, num, sizeof(int));
        setNum<<< 1 , 1>>>( num );

        cudaError_t cuda_status = cudaDeviceSynchronize();
        if (cuda_status != cudaSuccess) {
                printf("No dice\n");
                exit(1);
        }
        int results = 0;
        cudaMemcpyFromSymbol(&results, x, sizeof(int));
        printf("Value of x: %i\n", results);

        start = clock();
        for (int i = 0; i < 8; i++)
                Pressure<<<65535, 1024>>>(i);
        cuda_status = cudaDeviceSynchronize();
        printf("Result: %f\n", (float)(clock()-start)/CLOCKS_PER_SEC);
        return 0;
}

编译：nvcc -m64 -gencode arch=compute_20,code=sm_20 -o test test.cu

运行：./test 12（12 设置 x 变量）

注意注释掉的代码块：

跑步other(x, a);，我得到1.370000

跑步other(12,a);，我得到0.020000

4

2 回答 2

4

你还没有展示你的代码，所以我的评论本质上是一般性的。

通常，当我们看到基于对代码的相对较小更改而导致执行时间的巨大差异时，这是由于编译器可以优化的内容发生了变化。所以有几点需要考虑：

  1. 在我看来，用变量替换常量会导致这种级别的执行时间变化的想法似乎不太可能，因为编译器有很多方法可以优化对常用数据的访问，即使它本质上是动态/可变的。您可能想要比较每种情况下生成的 PTX 代码，以了解为什么存在这种差异，并测试您的结论，即实际差异是由于重复访问造成的。通常编译器会检测到这一点（特别是对于未修改的值）并优化对寄存器的访问。
  2. 如果命令行选项的数量相对较少，您可以考虑使用模板化内核，每个选项/选项都有不同的实例化。这应该会导致内核有效地硬编码选择，因此它的性能应该大致相当于您更快的情况。

编辑：既然你现在已经发布了一些代码，我会做一些额外的评论。

  1. 您的程序中有错误，您没有正确捕获。请进行适当的 cuda 错误检查以避免由此造成的任何混淆。一个错误是您使用cudaMemset符号__device__。

  2. 您发布的代码中的差异是由于编译器优化造成的。我不打算对此进行大量分析，因为您发布的代码似乎基本上是无意义的代码。但是有两种方法我可以支持这个断言。

    • -G使用开关编译代码。两种情况之间的时间变得相同（对我来说大约是 7 秒）。这将关闭所有编译器优化。在没有优化的情况下，代码的执行时间基本相同。

    • 查看 PTX 输出。在这两种情况下，使用-ptx开关编译您的代码。在“快速”的情况下，在 PTX 文件的末尾，我看到了这个全局函数定义：

      .visible .entry _Z8Pressurei(
    .param .u32 _Z8Pressurei_param_0
)
{



.loc 2 52 2
ret;
}

这段代码什么也不做。它只是一个带有 return 语句的空函数。编译器完全优化了函数行为。

在“慢”的情况下，压力函数中有大约 50 行实际代码。（而且整个 ptx 文件要大得多。）

于 2013-07-23T16:58:45.080 回答
0

您可以尝试添加__constant__到变量，即

__device__ __constant__ int a;

或者，一个块的线程之一（例如：0）可以将变量复制到__shared__内核中的一个变量中。

于 2013-07-23T16:57:00.193 回答