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 对象的选项，因此它必须为您的向量选择一个或另一个（主机或设备、CPU 或 GPU）。无论选择哪一个，如果在错误的环境中使用，向量可能（将会）中断。

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

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

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

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

因此，作为这个问题的答案：

那么，在CUDA中如何通过索引来实现功能呢？

我建议暂时将主机代码中的按索引功能与设备代码中的按索引功能分开（例如，两个单独的容器），对于设备代码中的按索引功能，您可以使用任何技术（其中不要使用或依赖 lambdas）在其他问题中涉及，例如 this one。

这是一个工作示例（我认为）演示了上面的注释，我们可以创建一个 lambda“类型”的向量，并将该向量的结果元素用作主机和设备代码中的 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
$

如上所述，此代码不是合理的代码。它是用来证明一个特定点的。