时间戳

首页 > 解决方案 > 在 CUDA 扩展 lambda 中捕获变量时出现内存错误

c++ - 在 CUDA 扩展 lambda 中捕获变量时出现内存错误

问题描述

我__device__在 CUDA 中创建了一个扩展的（即）lambda（参见例如这里），它应该捕获一个变量（这里是一个简单的double value = 3;）。它编译，但运行它，我得到一个invalid memory access错误，我不明白为什么。

更改变量以static const double value = 3解决问题，因为它不再被捕获（尽管我不明白它如何在 lambda 中仍然可用）。

问题 1：如何正确捕获 CUDA 扩展 lambda 中的主机变量？

问题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;
}

标签： c++c++11lambdacuda

解决方案

您需要按值传递 lambda，因为当您通过引用传递 lambda 时，lambda 中按值捕获的变量将在设备中不可用。

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

如果按值传递 lambda，则对象（及其内部）将被复制到内核。

推荐阅读