我已将我的项目缩减为仅包含相关代码。真正困扰我的部分是这不会产生任何错误。 无论如何,我有一个结构 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 内存变得不可变?我认为这是某种错误,但它编译和运行时没有警告或错误。
现在按照 promise ,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,你必须适本地修饰它,就像你打算在设备上执行的任何其他代码一样。初步介绍了这个概念 here ,尽管您可以找到许多其他示例。
根据您的最终意图,修复代码的方法可能有很多,但是与上述介绍/链接密切相关的方法可能如下所示:
$ 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__
装饰器是必需的。)
请注意,我正在处理 the original code you posted ,而不是@einpoklum 编辑到你的问题中的修改版本
在寻求他人帮助之前,如果您在使用 CUDA 代码时遇到问题,我通常建议您一定要这样做 proper CUDA error checking并使用 cuda-memcheck
运行您的代码。即使您不理解输出,它也会对那些试图帮助您的人有用。
关于c++ - 为什么将函数传递给内核会导致数据变得不可变?,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/51953435/