从结构启动 Cuda 调用

Launching Cuda Call from struct

本文关键字：调用 Cuda 启动 结构      更新时间：2023-10-16

给定一个简单的结构来包装 cuda 代码，可以写出类似的东西

func<float> s;
s.val = 3.f;
start_correct<<<1, 2>>>(s);

但是，我想将块、网格、共享内存计算放入结构中并调用内核，例如

func<float> s;
s.val = 3.f;
s.launch();

当第一个工作时，第二个给了我一个非法的内存访问错误。

重现我的问题的一个最小示例是

#include <stdio.h>
template<typename T>
struct func;
template<typename T>
__global__ void start(const func<T>& s){
printf("host access val %f n",s.val);
s();
}
template<typename T>
struct func
{
T val;
__device__ void operator()() const{
printf("device access val %f [%d]n",val,threadIdx.x);
}
enum{ C_N = 2 };
void launch()
{
start<<<1, C_N>>>(*this);
}
};
template<typename T>
__global__ void start_correct(const func<T> s){
printf("host access val %f n", s.val);
s();
}
int main(int argc, char const *argv[])
{
cudaError_t err;
func<float> s;
s.val = 3.f;
// launch cuda kernel <-- WORKS
start_correct<<<1, 2>>>(s);
cudaDeviceSynchronize();
if (err != cudaSuccess) printf("Error: %sn", cudaGetErrorString(err));

// launch cuda kernel <-- DOES NOT WORK
s.launch();
cudaDeviceSynchronize();
err = cudaGetLastError();
if (err != cudaSuccess) printf("Error: %sn", cudaGetErrorString(err));

return 0;
}

输出为

host access val 3.000000 
host access val 3.000000 
device access val 3.000000 [0]
device access val 3.000000 [1]
host access val 0.000000 
host access val 0.000000 
device access val 0.000000 [0]
device access val 0.000000 [1]
Error: an illegal memory access was encountered

两种方式不应该是等价的吗？是否有任何替代方案，也可以在结构内部进行shm，网格计算？

除非使用的是托管内存(不是)，否则通过引用传递内核参数是不合法的：

__global__ void start(const func<T>& s){
^

当我删除该与号时，您的代码运行时没有任何运行时错误，并给出了合理的输出：

$ cuda-memcheck ./t355
========= CUDA-MEMCHECK
host access val 3.000000
host access val 3.000000
device access val 3.000000 [0]
device access val 3.000000 [1]
host access val 3.000000
host access val 3.000000
device access val 3.000000 [0]
device access val 3.000000 [1]
========= ERROR SUMMARY: 0 errors
$

请注意，这实际上没有意义：

cudaDeviceSynchronize();
if (err != cudaSuccess) printf("Error: %sn", cudaGetErrorString(err));

并为我抛出编译器警告。

也许你的意思是：

err = cudaDeviceSynchronize();
if (err != cudaSuccess) printf("Error: %sn", cudaGetErrorString(err));
相关文章：