CUDA 7.5 thử nghiệm __host__ __device__ lambdas

Question

Tôi đã chơi một chút với số experimental device lambdas được giới thiệu trong CUDA 7.5 và được quảng cáo trong số blog post by Mark Harris này.

Ví dụ sau tôi đã xóa rất nhiều nội dung không cần thiết để hiển thị sự cố của mình (triển khai thực tế của tôi trông đẹp hơn một chút ...).

Tôi đã cố gắng viết một chức năng foreach hoạt động trên vectơ trên thiết bị (1 chuỗi cho mỗi phần tử) hoặc máy chủ (nối tiếp) tùy thuộc vào thông số mẫu. Với chức năng foreach này tôi có thể dễ dàng thực hiện các chức năng BLAS. Như một ví dụ tôi sử dụng gán một vô hướng để mỗi thành phần của một vector (tôi đính kèm mã hoàn chỉnh cuối cùng):

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); 
    } 
} 

Tuy nhiên, mã này đưa ra một lỗi biên dịch vì __host__ __device__ lambda:

kiểu đóng cửa cho một lambda ("lambda -> void") không thể được sử dụng trong các kiểu mẫu đối số của một hàm __global__ mẫu instantiation, trừ trường hợp lambda được định nghĩa trong một __device__ hoặc __global__ chức năng

tôi nhận được các lỗi tương tự nếu tôi loại bỏ các __device__ từ biểu thức lambda và tôi nhận được không có lỗi biên dịch nếu tôi loại bỏ __host__ (chỉ __device__ lambda), nhưng trong trường hợp này phần máy chủ không được thực hiện ...

Nếu tôi xác định lambda là hoặc là __host__ hoặc __device__ riêng biệt, mã biên dịch và hoạt động như mong đợi.

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); 
    } 
} 

Tuy nhiên, điều này giới thiệu mã trùng lặp và thực sự làm cho toàn bộ ý tưởng sử dụng lambdas vô dụng cho ví dụ này.

Có cách nào để hoàn thành những gì tôi muốn làm hay đây có phải là lỗi trong tính năng thử nghiệm không? Trên thực tế, việc xác định một lambda __host__ __device__ được đề cập rõ ràng trong ví dụ đầu tiên trong programming guide. Ngay cả đối với ví dụ đơn giản hơn (chỉ cần trả về một giá trị không đổi từ lambda), tôi không thể tìm thấy cách sử dụng biểu thức lambda trên cả máy chủ và thiết bị.

Dưới đây là toàn bộ mã, biên dịch với các tùy chọn -std=c++11 --expt-extended-lambda:

#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); 
    } 
} 

tôi đã sử dụng phiên bản sản xuất của CUDA 7.5.

Cập nhật

tôi đã cố gắng phiên bản thứ ba này cho hàm assignScalar:

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); 
    } 
} 

Nó biên dịch và chạy mà không có lỗi, nhưng phiên bản thiết bị (assignScalar3<true>) không được thực thi. Trên thực tế, tôi nghĩ rằng __CUDA_ARCH__ sẽ luôn luôn được undefined (kể từ khi chức năng không phải là __device__) nhưng tôi đã kiểm tra một cách rõ ràng rằng có một con đường biên dịch, nơi nó được xác định.

Answer 1

Nhiệm vụ mà tôi đã cố thực hiện với các ví dụ được cung cấp trong câu hỏi là không thể với CUDA 7.5, mặc dù nó không được loại trừ rõ ràng khỏi các trường hợp được phép cho hỗ trợ lambda thử nghiệm.

NVIDIA thông báo rằng CUDA Toolkit 8.0 sẽ hỗ trợ __host__ __device__ lambdas làm tính năng thử nghiệm, theo bài đăng trên blog CUDA 8 Features Revealed.

Tôi đã xác minh rằng ví dụ của tôi hoạt động với CUDA 8 Release Candidate (công cụ biên dịch Cuda, phát hành 8.0, V8.0.26).

Đây là mã mà cuối cùng tôi đã sử dụng, biên soạn với 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); 
    } 
}