CUDA C++ 11，lambdas数组，按索引功能，不起作用

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以来，这是一个无趣的演示。我们可以创建一个多元素向量，但是我们可以存储的唯一元素是您的两个lambda中的一个（不能同时存在）。

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

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

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

那么，我们怎样才能通过指数CUDA实现的功能？

我建议的时间由在主机代码索引作为该函数可以由函数由指数在设备代码分离（例如，两个单独的容器中），以及用于通过在设备代码索引功能，可以使用任何的技术（不使用或不依赖lambda）在其他问题中涵盖，如this one。

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

如上面已经提到的，该代码是不明智的代码。它是先进的证明一个特殊点。