CUDA學習（七十九）

關於__host____device__lambdas的說明:

與__device__ lambdas不同，可以從主機**呼叫__host____device__lambdas。如前所述，cuda編譯器用乙個已命名的佔位符型別的例項替換主機**中定義的擴充套件lambda表示式。擴充套件__host____device__lambda的佔位符型別通過間接函式呼叫呼叫orignal lambda的operator（）。

間接函式呼叫的存在可能會導致擴充套件的__host____device__lambda被主機編譯器優化得比只隱式或明確__host__的lambda更優化。在後一種情況下，主機編譯器可以輕鬆地將lambda的主體內聯到呼叫上下文中。但是如果擴充套件了__host____device__lambda，主機編譯器會遇到間接函式呼叫，並且可能無法輕鬆內聯原始__host____device__lambda主體。

*this capture by value:

當在非靜態類成員函式中定義lambda時，並且lambda的主體引用類成員變數時，c ++ 11 / c ++ 14規則要求該類的該指標由值捕獲，而不是被引用的成員變數。如果lambda是在主機函式中定義的擴充套件的__device__或__host____device__lambda，並且在gpu上執行lambda，則如果此指標指向主機記憶體，則在gpu上訪問引用的成員變數將導致執行時錯誤.

例子：

#include template __global__ void foo(t in) 
struct s1_t ;
void doit(void) ;
// kernel launch fails at run time because 'this->***'
// is not accessible from the gpu
foo << <1, 1 >> >(lam1);
cudadevicesynchronize();
}};int main(void)

c ++ 17通過新增新的「this」捕獲模式解決了這個問題。在這種模式下，編譯器複製由「 this」表示的物件，而不是按值捕獲指標。下面詳細描述「* this」捕捉模式：

當使用--expt-extended-lambda nvcc標誌時，cuda編譯器支援在__device__和__global__函式中定義的lambda以及在主機**中定義的擴充套件__device__lambda的「* this」捕獲模式。

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

#include template __global__ void foo(t in) 
struct s1_t ;
void doit(void) ;
// kernel launch succeeds
foo << <1, 1 >> >(lam1);
cudadevicesynchronize();
}};int main(void)

「*this」捕獲模式不允許用於在主機**中定義的未注釋的lambdas或擴充套件的__host__devicelambdas。支援和不支援的用法示例：

struct s1_t ;
void host_func(void) ;
// error: use in an extended __host__ __device__ lambda
auto lam2 = [=, *this] __host__ __device__;
// error: use in an unannotated lambda in host function
auto lam3 = [=, *this];
}__device__ void device_func(void) ;
// ok: use in a lambda defined in a __device__ function
auto lam2 = [=, *this] __host__ __device__;
// ok: use in a lambda defined in a __device__ function
auto lam3 = [=, *this];
}__host__ __device__ void host_device_func(void) ;
// error: use in an extended __host__ __device__ lambda
auto lam2 = [=, *this] __host__ __device__;
// error: use in an unannotated lambda in a __host__ __device__ function
auto lam3 = [=, *this];}};

補充筆記：

adl lookup：如前所述，在呼叫主機編譯器之前，cuda編譯器將用佔位符型別的例項替換擴充套件的lambda表示式。佔位符型別的乙個模板引數使用包含原始lambda表示式的函式的位址。對於引數型別涉及擴充套件lambda表示式的閉包型別的任何主機函式呼叫，這可能會導致其他命名空間參與引數相關查詢（adl）。這可能會導致主機編譯器選擇錯誤的功能。

例子：

namespace n1 ;
template void foo(t);
};namespace n2 
}void bar(n1::s1_t in)
;    n2::doit(lam1);
}

在上面的例子中，cuda編譯器用包含n1命名空間的佔位符型別替換了擴充套件lambda。因此，命名空間n1參與了n2 :: doit主體中的foo（in）的adl查詢，並且主機編譯失敗，因為找到了多個過載候選n1 :: foo和n2 :: foo。

CUDA學習（七十九）

六十九到七十九題

Spark七十九 Spark RDD API一

第七十九題合集微軟面試100題第七十九題

CUDA學習（七十九）

六十九到七十九題

Spark七十九 Spark RDD API一

第七十九題合集 微軟面試100題 第七十九題

相關推薦

第七十九題合集微軟面試100題第七十九題