为什么将函数传递给内核会导致数据变得不可变?
Why does passing a function to a kernel cause data to become immutable?
我已经将我的项目简化为仅相关的代码。真正困扰我的部分是这不会产生任何错误。 无论如何，我有一个结构 GpuData
struct GpuData { float x, y, z; };
我的目标是针对这个结构启动一个内核，该结构接受一个函数并将该函数应用于该结构。 因此，让我们看一个示例内核：
__global__ void StructFunctor(GpuData* in_dat, nvstd::function<float(void)> func) {
in_dat->x = func();
in_dat->y += T{1};
};
在这种情况下，内核被简化为非常简单的东西。它将 x 值设置为函数的结果。然后，它将在 y 值上加 1。
所以让我们试试吧。完整的源文件 (cuda_demo.cu(：
#include <iostream>
#include <nvfunctional>
struct GpuData { float x, y, z; };
__global__ void StructFunctor(GpuData* in_dat, nvstd::function<float(void)> func) {
in_dat->x = func();
in_dat->y += float{1};
};
int main(int argc, char** argv) {
GpuData c_dat {2, 3, 5};
std::cout << "Input x: " << c_dat.x << " y: " << c_dat.y << " z: " << c_dat.z << std::endl;
GpuData* g_dat;
cudaMalloc(&g_dat, sizeof(GpuData));
cudaMemcpy(g_dat, &c_dat, sizeof(GpuData), cudaMemcpyHostToDevice);
StructFunctor<<<1, 1>>>(g_dat, []()->float{return 1.0f;});
cudaMemcpy(&c_dat, g_dat, sizeof(GpuData), cudaMemcpyDeviceToHost);
std::cout << "Output x: " << c_dat.x << " y: " << c_dat.y << " z: " << c_dat.z << std::endl;
return 0;
}
好吧，如果我们真的要尝试它，我们将需要Cmake文件。我已经在最后贴上了这些。
在我的机器上，它可以编译和运行，没有错误。这是我的输出：
./CudaDemo
Input x: 2 y: 3 z: 5
Output x: 2 y: 3 z: 5
它们的变量根本没有被修改！但是如果我返回并注释掉in_dat-> = func();那么我会得到以下输出：
./CudaDemo
Input x: 2 y: 3 z: 5
Output x: 2 y: 4 z: 5
现在 y 值已被修改！这是一个好的开始，但是为什么当我尝试使用该函数时，GPU内存变得不可变？我认为这是某种错误，但它在没有警告或错误的情况下编译和运行。
现在正如承诺的那样，cmake 文件来运行它。
cmake_minimum_required(VERSION 3.8)
project(Temp LANGUAGES CXX CUDA)
set(CMAKE_CUDA_STANDARD 14)
add_executable(CudaDemo cuda_demo.cu)
set_property(TARGET CudaDemo PROPERTY CUDA_SEPARABLE_COMPILATION ON)
问题是您的代码正在主机代码中创建 lambda(因此它是针对您指定的任何主机处理器编译的(，然后您尝试在设备代码中使用该编译的 lambda。 这是行不通的。 如果您使用cuda-memcheck运行代码，则表示可能采取多种形式之一的错误，我看到"无效PC"的消息，这意味着您的程序试图从无效位置执行指令：
$ cuda-memcheck ./t277
========= CUDA-MEMCHECK
Input x: 2 y: 3 z: 5
========= Invalid PC
========= at 0x00000048 in void StructFunctor<float>(GpuData<float>*, nvstd::function<float () (void)>)
========= by thread (0,0,0) in block (0,0,0)
========= Device Frame:void StructFunctor<float>(GpuData<float>*, nvstd::function<float () (void)>) (void StructFunctor<float>(GpuData<float>*, nvstd::function<float () (void)>) : 0x40)
========= Saved host backtrace up to driver entry point at kernel launch time
========= Host Frame:/lib64/libcuda.so.1 (cuLaunchKernel + 0x2cd) [0x2486ed]
========= Host Frame:./t277 [0x190b2]
========= Host Frame:./t277 [0x192a7]
在 CUDA 中，如果要在设备代码中使用 lambda，则必须正确修饰它，就像您打算在设备上执行的任何其他代码一样。 这里初步介绍了这个概念，尽管你可以找到许多其他例子。
可能有很多方法可以修复代码，具体取决于您的最终意图，但是与上述介绍/链接密切相关的方法可能如下所示：
$ cat t277.cu
#include <iostream>
template <typename T>
struct GpuData {
T x;
T y;
T z;
};
template <typename T, typename F>
__global__ void StructFunctor(GpuData<T>* in_dat, F f) {
in_dat->x = f();
in_dat->y += T{1};
};
int main(int argc, char** argv) {
GpuData<float> c_dat {2, 3, 5};
std::cout << "Input x: " << c_dat.x << " y: " << c_dat.y << " z: " << c_dat.z << std::endl;
GpuData<float>* g_dat;
cudaMalloc(&g_dat, sizeof(GpuData<float>));
cudaMemcpy(g_dat, &c_dat, sizeof(GpuData<float>), cudaMemcpyHostToDevice);
StructFunctor<float><<<1, 1>>>(g_dat, [] __host__ __device__ ()->float{return 1.0f;});
cudaMemcpy(&c_dat, g_dat, sizeof(GpuData<float>), cudaMemcpyDeviceToHost);
std::cout << "Output x: " << c_dat.x << " y: " << c_dat.y << " z: " << c_dat.z << std::endl;
return 0;
}
$ nvcc -std=c++11 t277.cu -o t277 --expt-extended-lambda
$ cuda-memcheck ./t277
========= CUDA-MEMCHECK
Input x: 2 y: 3 z: 5
Output x: 1 y: 4 z: 5
========= ERROR SUMMARY: 0 errors
$
(在这种特殊情况下，我添加到 lambda 的__host__装饰器不是必需的，但__device__装饰器是必需的。
请注意，我正在处理您发布的原始代码，而不是由@einpoklum编辑到您的问题中的修改版本
在向他人寻求帮助之前，如果您在使用 CUDA 代码时遇到问题，我通常建议您确保进行正确的 CUDA 错误检查并使用cuda-memcheck运行您的代码。 即使您不了解输出，它也会对那些试图帮助您的人有用。
事实上，正如@RobertCrovella所指出的，(仅主机(lambda 的地址不是有效的设备端地址，因此构造的 nvstd：：function 不可调用(。当您尝试在内核中调用它时，会出现错误。这是您的代码(好吧，我对代码的编辑(，转换为使用正确的错误检查：
#include <iostream>
#include <nvfunctional>
#include <cuda/api_wrappers.h>
struct GpuData { float x, y, z; };
__global__ void StructFunctor(GpuData* in_dat, nvstd::function<float(void)> func) {
in_dat->x = func();
in_dat->y += float{1};
};
int main(int argc, char** argv) {
using std::cout; using std::endl;
GpuData c_dat {2, 3, 5};
cout << "Input x: " << c_dat.x << " y: " << c_dat.y << " z: " << c_dat.z << endl;
auto current_device = cuda::device::current::get();
auto g_dat = cuda::memory::device::make_unique<GpuData>(current_device);
cuda::memory::copy(g_dat.get(), &c_dat, sizeof(GpuData));
device.launch(StructFunctor, cuda::make_launch_config(1, 1),
g_dat.get(), []()->float { return 1.0f; });
cuda::outstanding_error::ensure_none(); // This is where we'll notice the error
cuda::memory::copy(&c_dat, g_dat.get(), sizeof(GpuData));
cout << "Output x: " << c_dat.x << " y: " << c_dat.y << " z: " << c_dat.z << std::endl;
}
运行此命令时，您将获得：
Input x: 2 y: 3 z: 5
terminate called after throwing an instance of 'cuda::runtime_error'
what(): Synchronously copying data: an illegal memory access was encountered
Aborted
解决方法是：
#include <iostream>
#include <cuda/api_wrappers.h>
struct GpuData { float x, y, z; };
template <typename F>
__global__ void StructFunctor(GpuData* in_dat, F func) {
in_dat->x = func();
in_dat->y += float{1};
};
int main(int argc, char** argv) {
using std::cout; using std::endl;
GpuData c_dat {2, 3, 5};
cout << "Input x: " << c_dat.x << " y: " << c_dat.y << " z: " << c_dat.z << endl;
auto device = cuda::device::current::get();
auto g_dat = cuda::memory::device::make_unique<GpuData>(device);
cuda::memory::copy(g_dat.get(), &c_dat, sizeof(GpuData));
auto return_one = [] __device__ ()->float { return 1.0f; };
device.launch(StructFunctor<decltype(return_one)>, cuda::make_launch_config(1, 1), g_dat.get(), return_one);
cuda::outstanding_error::ensure_none();
cuda::memory::copy(&c_dat, g_dat.get(), sizeof(GpuData));
cout << "Output x: " << c_dat.x << " y: " << c_dat.y << " z: " << c_dat.z << endl;
}
要使用 CUDA API 包装器，请将其添加到您的CMakeLists.txt
ExternalProject_Add(cuda-api-wrappers_project
PREFIX CMakeFiles/cuda-api-wrappers_project
TMP_DIR CMakeFiles/cuda-api-wrappers_project/tmp
STAMP_DIR CMakeFiles/cuda-api-wrappers_project/stamp
GIT_REPOSITORY git@github.com:eyalroz/cuda-api-wrappers.git
GIT_TAG 7e48712af95939361bf04e4f4718688795a319f9
UPDATE_COMMAND ""
SOURCE_DIR "${CMAKE_SOURCE_DIR}/cuda-api-wrappers"
BUILD_IN_SOURCE 1
INSTALL_COMMAND ""
)
相关文章：
- 正在寻找C++不可变的hashset/hashmap
- 具有常量属性的不可变类
- 为什么将函数传递给内核会导致数据变得不可变?
- 使用 Broadcast 发出的从节点服务器发送的数据不能被 C++ 套接字 IO 客户端读取
- 现代c++编译器会优化不可变的临时变量吗
- Winsock 数据不是从 IP 检索的，而是从普通 URL 检索的
- C++覆盖父结构中的数据不起作用
- 使用ifstream时，被调用文件中的数据不会打印
- C++中的不可变列表
- C++/WinRT 数据 不正确的数据封送处理
- 如何将地图作为不可变地图传递?
- 数据不是用 ofstream 写入的，即使返回成功也是如此
- 为什么数据不传输
- 不可变的全局对象应该声明为"const my_result_t BLAH"还是"extern const my_result_t BLAH;"？
- 向不可变纹理存储提供数据
- 在C++中创建不可变且高效的类的惯用方法
- std：：shared_ptr<std：：string const> 可以作为引用计数的不可变字符串的有效实现吗？
- 不可变数据模型的内存管理
- 如何使用分支逻辑构造不可变数据类型的实例
- C++11 中的不可变"functional"数据结构
