CUDA學習（七十九）

關于<code>__host__</code> <code>__device__</code> lambdas的說明:

與<code>__device__ lambdas</code>不同，可以從主機代碼調用<code>__host__</code> <code>__device__</code> lambdas。 如前所述，CUDA編譯器用一個已命名的占位符類型的執行個體替換主機代碼中定義的擴充lambda表達式。 擴充<code>__host__</code> <code>__device__</code>lambda的占位符類型通過間接函數調用調用orignal lambda的operator（）。

間接函數調用的存在可能會導緻擴充的<code>__host__</code> <code>__device__</code> lambda被主機編譯器優化得比隻隐式或明确<code>__host__</code>的lambda更優化。 在後一種情況下，主機編譯器可以輕松地将lambda的主體内聯到調用上下文中。 但是如果擴充了<code>__host__</code> <code>__device__</code> lambda，主機編譯器會遇到間接函數調用，并且可能無法輕松内聯原始<code>__host__</code> <code>__device__</code> lambda主體。

*this Capture By Value:

當在非靜态類成員函數中定義lambda時，并且lambda的主體引用類成員變量時，C ++ 11 / C ++ 14規則要求該類的該指針由值捕獲， 而不是被引用的成員變量。 如果lambda是在主機函數中定義的擴充的<code>__device__</code>或<code>__host__</code> <code>__device__</code>lambda，并且在GPU上執行lambda，則如果此指針指向主機記憶體，則在GPU上通路引用的成員變量将導緻運作時錯誤.

例子：

當使用--expt-extended-lambda nvcc标志時，CUDA編譯器支援在<code>__device__</code>和<code>__global__</code>函數中定義的lambda以及在主機代碼中定義的擴充<code>__device__</code> lambda的“* this”捕獲模式。

上面的示例修改為使用“* this”捕獲模式：

“*this”捕獲模式不允許用于在主機代碼中定義的未注釋的lambdas或擴充的__host__ device lambdas。 支援和不支援的用法示例：

補充筆記：

ADL Lookup：如前所述，在調用主機編譯器之前，CUDA編譯器将用占位符類型的執行個體替換擴充的lambda表達式。 占位符類型的一個模闆參數使用包含原始lambda表達式的函數的位址。 對于參數類型涉及擴充lambda表達式的閉包類型的任何主機函數調用，這可能會導緻其他名稱空間參與參數相關查找（ADL）。 這可能會導緻主機編譯器選擇錯誤的功能。

在上面的例子中，CUDA編譯器用包含N1命名空間的占位符類型替換了擴充lambda。 是以，名稱空間N1參與了N2 :: doit主體中的foo（in）的ADL查找，并且主機編譯失敗，因為找到了多個過載候選N1 :: foo和N2 :: foo。