c++ - CUDA C++11，lambda 数组，按索引函数，不工作

Question

我在尝试让 CUDA 程序通过索引管理 lambda 数组时遇到了问题。重现问题的示例代码

 #include <cuda.h>
 #include <vector>
 #include <stdio.h>
 #include <stdlib.h>
 #include <time.h>
 #include <sys/time.h>
 #include <cassert>

 #define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
 inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=true){
     if (code != cudaSuccess) {
         fprintf(stderr,"GPUassert: %s %s %d\n",
         cudaGetErrorString(code), file, line);
         if (abort) exit(code);
     }   
 }

 template<typename Lambda>
 __global__ void kernel(Lambda f){ 
     int t = blockIdx.x * blockDim.x + threadIdx.x;
     printf("device: thread %i: ", t); 
     printf("f() = %i\n", f() );
 }

 int main(int argc, char **argv){
     // arguments
     if(argc != 2){ 
         fprintf(stderr, "run as ./prog i\nwhere 'i' is function index");
         exit(EXIT_FAILURE);
     }   
     int i = atoi(argv[1]);


     // lambdas
     auto lam0 = [] __host__ __device__ (){ return 333; };
     auto lam1 = [] __host__ __device__ (){ return 777; };


     // make vector of functions
     std::vector<int(*)()> v;
     v.push_back(lam0);
     v.push_back(lam1);


     // host: calling a function by index
     printf("host: f() = %i\n", (*v[i])() );


     // device: calling a function by index
     kernel<<< 1, 1 >>>( v[i] ); // does not work
     //kernel<<< 1, 1 >>>( lam0 ); // does work
     gpuErrchk( cudaPeekAtLastError() );
     gpuErrchk( cudaDeviceSynchronize() );
     return EXIT_SUCCESS;
 }

编译

nvcc -arch sm_60 -std=c++11 --expt-extended-lambda main.cu -o prog

运行时出现的错误是

➜  cuda-lambda ./prog 0
host: f() = 333
device: GPUassert: invalid program counter main.cu 53

似乎 CUDA 无法管理 int(*)() 函数指针形式(而宿主 C++ 确实可以正常工作)。另一方面，每个 lambda 都作为不同的数据类型进行管理，无论它们在代码中是否相同并且具有相同的契约。那么，在CUDA中如何通过索引实现功能呢？

Answer 1

这里有一些注意事项。

尽管您建议要“管理 lambda 数组”，但您实际上依赖于将 lambda 优雅地转换为函数指针(当 lambda 未捕获时可能)。

当您将某些东西标记为 __host__ __device__ 时，您是在向编译器声明需要编译该项目的两个拷贝(具有两个明显不同的入口点):一个用于 CPU，另一个用于 GPU。

当我们采用 __host__ __device__ lambda 并要求它降级为函数指针时，我们会面临“选择哪个函数指针(入口点)？”的问题。编译器不再具有携带实验性 lambda 对象的选项，因此它必须为您的 vector 选择一个或另一个(主机或设备、CPU 或 GPU)。无论选择哪一个，如果在错误的环境中使用， vector 都可能(将)中断。

由此得出的一个结论是，您的两个测试用例不相同。在一种情况下(损坏)您将函数指针传递给内核(因此内核被模板化以接受函数指针参数)而在另一种情况下(工作)您将 lambda 传递给内核(因此内核被模板化接受 lambda 参数)。

在我看来，这里的问题不仅仅是由容器的使用引起的，而是由您使用的容器类型引起的。我可以通过将您的 vector 转换为实际 lambda 类型的 vector ，以一种简单的方式(见下文)来证明这一点。在那种情况下，我们可以使代码“工作”(某种程度上)，但由于 every lambda has a unique type ，这是一个无趣的演示。我们可以创建一个多元素 vector ，但我们可以在其中存储的唯一元素是您的两个 lambda 之一(不能同时)。

如果我们使用可以处理不同类型的容器(例如 std::tuple)，也许我们可以在这里取得一些进展，但我知道没有直接的方法来索引此类元素一个容器。即使我们可以，接受 lambda 作为参数/模板类型的模板内核也必须为每个 lambda 实例化。

在我看来，函数指针避免了这种特殊类型的“困惑”。

因此，作为对这个问题的回答:

Then, how can we achieve function by index in CUDA?

我建议暂时将主机代码中的索引函数与设备代码中的索引函数分开(例如，两个单独的容器)，并且对于设备代码中的索引函数，您可以使用任何技术(不要使用或依赖 lambdas) 涵盖在其他问题中，例如 this one .

这是一个有效的例子(我认为)证明了上面的注释，我们可以创建一个 lambda“类型”的 vector ，并在主机和设备代码中使用该 vector 的结果元素作为 lambda:

$ cat t64.cu
 #include <cuda.h>
 #include <vector>
 #include <stdio.h>
 #include <stdlib.h>
 #include <time.h>
 #include <sys/time.h>
 #include <cassert>

 #define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
 inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=true){
     if (code != cudaSuccess) {
         fprintf(stderr,"GPUassert: %s %s %d\n",
         cudaGetErrorString(code), file, line);
         if (abort) exit(code);
     }
 }


 template<typename Lambda>
 __global__ void kernel(Lambda f){
     int t = blockIdx.x * blockDim.x + threadIdx.x;
     printf("device: thread %i: ", t);
     printf("f() = %i\n", f() );
 }

 template <typename T>
 std::vector<T> fill(T L0, T L1){
   std::vector<T> v;
   v.push_back(L0);
   v.push_back(L1);
   return v;
}

 int main(int argc, char **argv){
     // arguments
     if(argc != 2){
         fprintf(stderr, "run as ./prog i\nwhere 'i' is function index");
         exit(EXIT_FAILURE);
     }
     int i = atoi(argv[1]);


     // lambdas
     auto lam0 = [] __host__ __device__ (){ return 333; };
     auto lam1 = [] __host__ __device__ (){ return 777; };

     auto v = fill(lam0, lam0);

     // make vector of functions
 //    std::vector< int(*)()> v;
 //    v.push_back(lam0);
 //    v.push_back(lam1);


     // host: calling a function by index
     // host: calling a function by index
     printf("host: f() = %i\n", (*v[i])() );


     // device: calling a function by index
     kernel<<< 1, 1 >>>( v[i] ); // does not work
     //kernel<<< 1, 1 >>>( lam0 ); // does work
     gpuErrchk( cudaPeekAtLastError() );
     gpuErrchk( cudaDeviceSynchronize() );
     return EXIT_SUCCESS;
 }

$ nvcc -arch sm_61 -std=c++11 --expt-extended-lambda t64.cu -o t64
$ cuda-memcheck ./t64 0
========= CUDA-MEMCHECK
host: f() = 333
device: thread 0: f() = 333
========= ERROR SUMMARY: 0 errors
$ cuda-memcheck ./t64 1
========= CUDA-MEMCHECK
host: f() = 333
device: thread 0: f() = 333
========= ERROR SUMMARY: 0 errors
$

如上所述，此代码不是合理的代码。证明一个特定的观点是先进的。