Ошибка памяти при захвате переменной в расширенной лямбде CUDA - PullRequest
Новая
       8

Ошибка памяти при захвате переменной в расширенной лямбде CUDA

0 голосов
/ 21 октября 2019

Я создаю расширенную (т.е. __device__) лямбду в CUDA (см., Например, здесь ), и предполагается, что она захватывает переменную (здесь простая double value = 3;). Он компилируется, но при запуске я получаю ошибку invalid memory access, и я не понимаю, почему.

Изменение переменной на static const double value = 3 устраняет проблему, поскольку она больше не перехватывается (хотя я неЯ не понимаю, как это все еще доступно внутри лямбды).

Вопрос1: Как правильно перехватить переменные хоста в расширенной лямбде CUDA?

Вопрос2: почему этот код не работает?

Я пробовал это на Ubuntu 16, как с CUDA 8, так и с 10.

Код MWE

Скомпилировано с nvcc mwe_lambda.cu -o mwe_lambda --std=c++11 -lineinfo -arch=sm_60 --expt-relaxed-constexpr --expt-extended-lambda

Обратите внимание, в частности, lambda, который должен быть скопирован. managed_allocator и т. Д. Только для того, чтобы использовать управляемую память и распечатать ошибку CUDA. 

#include <cuda.h>
#include <cuda_runtime.h>

#include <vector>
#include <iostream>
#include <string>


static void CudaHandleError( cudaError_t err, const char *file, int line, const std::string & function)
{
    if (err != cudaSuccess)
    {
        std::cerr << std::string(cudaGetErrorString( err )) << " " << file << " " << line << " " << function << std::endl;
    }
}

#define CU_HANDLE_ERROR( err ) (CudaHandleError( err, __FILE__, __LINE__, __func__ ))

#define CU_CHECK_ERROR( ) (CudaHandleError( cudaGetLastError(), __FILE__, __LINE__, __func__ ))

#define CU_CHECK_AND_SYNC( ) CU_CHECK_ERROR(); CU_HANDLE_ERROR( cudaDeviceSynchronize() )


template<class T>
class managed_allocator : public std::allocator<T>
{
public:
    using value_type = T;

    template<typename _Tp1>
    struct rebind
    {
        typedef managed_allocator<_Tp1> other;
    };

    value_type* allocate(size_t n)
    {
        value_type* result = nullptr;

        CU_HANDLE_ERROR( cudaMallocManaged(&result, n*sizeof(value_type)) );

        return result;
    }

    void deallocate(value_type* ptr, size_t)
    {
        CU_HANDLE_ERROR( cudaFree(ptr) );
    }

    managed_allocator() throw(): std::allocator<T>() { } //fprintf(stderr, "Hello managed allocator!\n"); }
    managed_allocator(const managed_allocator &a) throw(): std::allocator<T>(a) { }
    template <class U>                    
    managed_allocator(const managed_allocator<U> &a) throw(): std::allocator<T>(a) { }
    ~managed_allocator() throw() { }
};

template<typename T>
using field = std::vector<T, managed_allocator<T>>;

// vf[i] = f()
template<typename A, typename F>
__global__ void cu_set_lambda(A * vf, const F & f, int N)
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if(idx < N)
    {
        vf[idx] = f();
    }
}

int main()
{
    std::cerr << "started" << std::endl;
    {
        field<double> vf(10, 0);

        double value = 3;
        auto lambda = [=] __device__ ()
        {
            return value;
        };

        auto n = vf.size();
        cu_set_lambda<<<(n+1023)/1024, 1024>>>(vf.data(), lambda, n);
        CU_CHECK_AND_SYNC();

        std::cerr << vf[0] << " " << vf[1] << std::endl;
    }
    std::cerr << "finished" << std::endl;
}

1 Ответ

3 голосов
/ 21 октября 2019

Вам нужно передать лямбду по значению, так как переменные, захваченные по значению в лямбде, не будут доступны в устройстве, когда вы передадите лямбду по ссылке. 

__global__ void cu_set_lambda(A * vf, const F  f, int N)
                                      ^^^^^^^

Если вы передаете лямбду позначение, объект (и его внутренние компоненты) будут скопированы в ядро.

...