Cuda传递值更快？

Question

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

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

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

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

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

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

任何帮助将不胜感激。

*运行2.0硬件。

编辑：这是我遇到的不同的例子：./test 12（12套X变量）

注的注释：与nvcc -m64 -gencode arch=compute_20,code=sm_20 -o test test.cu

运行：编译时

#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; 
} 

out代码块：

正在运行other(x, a);，我得到1.370000

运行other(12,a);，我得到0.020000

Answer 1

你还没有显示你的代码，所以我的评论必然是一般性的。

通常，当我们根据代码的相对较小的变化看到执行时间差异很大时，这是由于编译器可以优化的内容的变化。因此，有几件事情要考虑：

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文件要大得多）。

Answer 2

你可以尝试添加__constant__的变量，即

__device__ __constant__ int a; 

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