सीयूडीए 7.5 प्रयोगात्मक __host__ __device__ lambdas

Question

मैंने experimental device lambdas के साथ थोड़ा सा खेला जहां सीयूडीए 7.5 में पेश किया गया और इस blog post by Mark Harris में प्रचारित किया गया।

निम्नलिखित उदाहरण के लिए मैंने अपनी सारी समस्याएं दिखाने के लिए बहुत सारी चीजें हटा दी हैं (मेरा वास्तविक कार्यान्वयन थोड़ा अच्छा लगता है ...)।

मैंने एक फ़ोरैच फ़ंक्शन लिखने की कोशिश की जो टेम्पलेट पैरामीटर के आधार पर डिवाइस पर वैक्टर (1 तत्व प्रति थ्रेड) या होस्ट (सीरियल) पर काम करता है। इस foreach समारोह के साथ मैं आसानी से बीएलएएस कार्यों को लागू कर सकते हैं। एक उदाहरण मैं एक वेक्टर के प्रत्येक घटक के लिए एक अदिश बताए उपयोग के रूप में (मैं अंत में पूरा कोड देते हैं):

template<bool onDevice> void assignScalar(size_t size, double* vector, double a) 
{ 
    auto assign = [=] __host__ __device__ (size_t index) { vector[index] = a; }; 
    if(onDevice) 
    { 
     foreachDevice(size, assign); 
    } 
    else 
    { 
     foreachHost(size, assign); 
    } 
} 

हालांकि, इस कोड __host__ __device__ लैम्ब्डा की वजह से एक संकलक त्रुटि देता है:

एक लैम्ब्डा ("लैम्ब्डा -> शून्य") के लिए बंद करने का प्रकार, एक __global__ समारोह टेम्पलेट इन्स्टेन्शियशन के टेम्पलेट तर्क प्रकार में नहीं किया जा सकता जब तक कि लैम्ब्डा एक __device__ या __global__ समारोह

मैं भीतर परिभाषित किया गया हैएक ही गलती करता है, तो मैं लैम्ब्डा अभिव्यक्ति से __device__ को हटा दें और अगर मैं __host__ (केवल __device__ लैम्ब्डा) को हटाने मुझे कोई संकलन त्रुटि मिलती है, लेकिन इस मामले में मेजबान हिस्सा निष्पादित नहीं है ...

अगर मैं लैम्ब्डा के रूप में परिभाषित या तो __host__ या __device__ अलग से, कोड संकलित और अपेक्षा के अनुसार काम करता है।

template<bool onDevice> void assignScalar2(size_t size, double* vector, double a) 
{ 
    if(onDevice) 
    { 
     auto assign = [=] __device__ (size_t index) { vector[index] = a; }; 
     foreachDevice(size, assign); 
    } 
    else 
    { 
     auto assign = [=] __host__ (size_t index) { vector[index] = a; }; 
     foreachHost(size, assign); 
    } 
} 

हालांकि, इस कोड दोहराव का परिचय है और वास्तव में lambdas इस उदाहरण के लिए बेकार का उपयोग कर के पूरे विचार बनाता है।

क्या मैं ऐसा करना चाहता हूं जो मैं करना चाहता हूं या यह प्रयोगात्मक सुविधा में एक बग है? दरअसल, __host__ __device__ लैम्ब्डा को परिभाषित करना programming guide में पहले उदाहरण में स्पष्ट रूप से उल्लेख किया गया है। यहां तक ​​कि उस सरल उदाहरण के लिए (केवल लैम्ब्डा से निरंतर मूल्य लौटाएं) मुझे मेजबान और डिवाइस दोनों पर लैम्ब्डा अभिव्यक्ति का उपयोग करने का कोई तरीका नहीं मिला।

#include <iostream> 
using namespace std; 

template<typename Operation> void foreachHost(size_t size, Operation o) 
{ 
    for(size_t i = 0; i < size; ++i) 
    { 
     o(i); 
    } 
} 

template<typename Operation> __global__ void kernel_foreach(Operation o) 
{ 
    size_t index = blockIdx.x * blockDim.x + threadIdx.x; 
    o(index); 
} 

template<typename Operation> void foreachDevice(size_t size, Operation o) 
{ 
    size_t blocksize = 32; 
    size_t gridsize = size/32; 
    kernel_foreach<<<gridsize,blocksize>>>(o); 
} 

__global__ void printFirstElementOnDevice(double* vector) 
{ 
    printf("dVector[0] = %f\n", vector[0]); 
} 

void assignScalarHost(size_t size, double* vector, double a) 
{ 
    auto assign = [=] (size_t index) { vector[index] = a; }; 
    foreachHost(size, assign); 
} 

void assignScalarDevice(size_t size, double* vector, double a) 
{ 
    auto assign = [=] __device__ (size_t index) { vector[index] = a; }; 
    foreachDevice(size, assign); 
} 

// compile error: 
template<bool onDevice> void assignScalar(size_t size, double* vector, double a) 
{ 
    auto assign = [=] __host__ __device__ (size_t index) { vector[index] = a; }; 
    if(onDevice) 
    { 
     foreachDevice(size, assign); 
    } 
    else 
    { 
     foreachHost(size, assign); 
    } 
} 

// works: 
template<bool onDevice> void assignScalar2(size_t size, double* vector, double a) 
{ 
    if(onDevice) 
    { 
     auto assign = [=] __device__ (size_t index) { vector[index] = a; }; 
     foreachDevice(size, assign); 
    } 
    else 
    { 
     auto assign = [=] __host__ (size_t index) { vector[index] = a; }; 
     foreachHost(size, assign); 
    } 
} 

int main() 
{ 
    size_t SIZE = 32; 

    double* hVector = new double[SIZE]; 
    double* dVector; 
    cudaMalloc(&dVector, SIZE*sizeof(double)); 

    // clear memory 
    for(size_t i = 0; i < SIZE; ++i) 
    { 
     hVector[i] = 0; 
    } 
    cudaMemcpy(dVector, hVector, SIZE*sizeof(double), cudaMemcpyHostToDevice); 

    assignScalarHost(SIZE, hVector, 1.0); 
    cout << "hVector[0] = " << hVector[0] << endl; 

    assignScalarDevice(SIZE, dVector, 2.0); 
    printFirstElementOnDevice<<<1,1>>>(dVector); 
    cudaDeviceSynchronize(); 

    assignScalar2<false>(SIZE, hVector, 3.0); 
    cout << "hVector[0] = " << hVector[0] << endl; 

    assignScalar2<true>(SIZE, dVector, 4.0); 
    printFirstElementOnDevice<<<1,1>>>(dVector); 
    cudaDeviceSynchronize(); 

// assignScalar<false>(SIZE, hVector, 5.0); 
// cout << "hVector[0] = " << hVector[0] << endl; 
// 
// assignScalar<true>(SIZE, dVector, 6.0); 
// printFirstElementOnDevice<<<1,1>>>(dVector); 
// cudaDeviceSynchronize(); 

    cudaError_t error = cudaGetLastError(); 
    if(error!=cudaSuccess) 
    { 
     cout << "ERROR: " << cudaGetErrorString(error); 
    } 
} 

मैं CUDA 7.5 के उत्पादन रिहाई का प्रयोग किया:

यहां विकल्पों -std=c++11 --expt-extended-lambda साथ संकलन पूर्ण कोड है।

template<bool onDevice> void assignScalar3(size_t size, double* vector, double a) 
{ 
#ifdef __CUDA_ARCH__ 
#define LAMBDA_HOST_DEVICE __device__ 
#else 
#define LAMBDA_HOST_DEVICE __host__ 
#endif 

    auto assign = [=] LAMBDA_HOST_DEVICE (size_t index) { vector[index] = a; }; 
    if(onDevice) 
    { 
     foreachDevice(size, assign); 
    } 
    else 
    { 
     foreachHost(size, assign); 
    } 
} 

यह कम्पाइल और त्रुटि के बिना चलाता है, लेकिन डिवाइस संस्करण (assignScalar3<true>) निष्पादित नहीं है:

अद्यतन

मैं assignScalar समारोह के लिए इस तीसरे संस्करण की कोशिश की। असल में, मैंने सोचा कि __CUDA_ARCH__ हमेशा अपरिभाषित होगा (क्योंकि फ़ंक्शन __device__ नहीं है) लेकिन मैंने स्पष्ट रूप से जांच की है कि एक संकलित पथ है जहां इसे परिभाषित किया गया है। हालांकि यह स्पष्ट रूप से प्रयोगात्मक लैम्ब्डा समर्थन के लिए अनुमति दी मामलों से बाहर रखा नहीं किया गया था

Answer 1

काम है कि मैं प्रश्न में दिए गए उदाहरणों के साथ पूरा करने के लिए कोशिश की, CUDA 7.5 साथ नहीं संभव है।

NVIDIA घोषणा की कि CUDA टूलकिट 8.0, एक प्रायोगिक सुविधा के रूप में __host__ __device__ lambdas का समर्थन करेंगे ब्लॉग पोस्ट CUDA 8 Features Revealed के अनुसार।

मैंने सत्यापित किया कि मेरा उदाहरण CUDA 8 रिलीज उम्मीदवार (कुडा संकलन उपकरण, रिलीज 8.0, V8.0.26) के साथ काम करता है।

यहाँ कोड है कि मैं अंत में इस्तेमाल किया, nvcc -std=c++11 --expt-extended-lambda साथ संकलित है:

#include <iostream> 
using namespace std; 

template<typename Operation> __global__ void kernel_foreach(Operation o) 
{ 
    size_t i = blockIdx.x * blockDim.x + threadIdx.x; 
    o(i); 
} 

template<bool onDevice, typename Operation> void foreach(size_t size, Operation o) 
{ 
    if(onDevice) 
    { 
     size_t blocksize = 32; 
     size_t gridsize = size/32; 
     kernel_foreach<<<gridsize,blocksize>>>(o); 
    } 
    else 
    { 
     for(size_t i = 0; i < size; ++i) 
     { 
      o(i); 
     } 
    } 
} 

__global__ void printFirstElementOnDevice(double* vector) 
{ 
    printf("dVector[0] = %f\n", vector[0]); 
} 

template<bool onDevice> void assignScalar(size_t size, double* vector, double a) 
{ 
    auto assign = [=] __host__ __device__ (size_t i) { vector[i] = a; }; 
    foreach<onDevice>(size, assign); 
} 

int main() 
{ 
    size_t SIZE = 32; 

    double* hVector = new double[SIZE]; 
    double* dVector; 
    cudaMalloc(&dVector, SIZE*sizeof(double)); 

    // clear memory 
    for(size_t i = 0; i < SIZE; ++i) 
    { 
     hVector[i] = 0; 
    } 
    cudaMemcpy(dVector, hVector, SIZE*sizeof(double), cudaMemcpyHostToDevice); 

    assignScalar<false>(SIZE, hVector, 3.0); 
    cout << "hVector[0] = " << hVector[0] << endl; 

    assignScalar<true>(SIZE, dVector, 4.0); 
    printFirstElementOnDevice<<<1,1>>>(dVector); 
    cudaDeviceSynchronize(); 

    cudaError_t error = cudaGetLastError(); 
    if(error!=cudaSuccess) 
    { 
     cout << "ERROR: " << cudaGetErrorString(error); 
    } 
}