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 
$ 

如上面已經提到的，該代碼是不明智的代碼。它是先進的證明一個特殊點。