Pycuda通过函数参数值声明数组大小时出现nvcc编译错误

1. 你用到的int d[in_integer]、int d[in_matrix[0]]属于变长数组（VLA），是C99标准引入的语法特性，而C标准本身没有原生支持该特性。nvcc作为CUDA C编译器，默认模式下不支持device端栈上的变长数组声明，因此会触发编译报错。而int d[1]是编译期可确定大小的固定长度数组，符合C++语法规则，因此可以正常编译运行。
2. 就算通过编译选项强行开启VLA支持，CUDA kernel的每线程栈空间非常有限（通常仅几KB），动态长度的栈数组极容易触发栈溢出，导致运行时异常，本身也不推荐这种写法。

PyCUDA本身只是对nvcc编译流程做了封装，是否支持该写法完全取决于nvcc的编译规则，你可以采用以下几种合规的方案实现可变长度数组的需求：

• 方案一：使用动态共享内存，这是CUDA官方推荐的可变长度数组实现方式。写法如下：

  __global__ void array_declaration(int *in_matrix, int *out_matrix, int in_integer)
  {
    extern __shared__ int d[]; // 动态共享内存声明
    out_matrix[0] = in_matrix[0];
  }
  

  调用kernel时传入对应大小的shared参数（单位为字节）：

  func(a_gpu, b_gpu, np.int32(1), block=(1,1,1), shared=1*4) # 1个int占4字节
  

• 方案二：如果数组长度在Host侧就可以提前确定，可以直接在Python侧将长度插值到kernel字符串中，生成固定长度的数组声明，完全避开VLA问题：

  array_len = 1
  mod = SourceModule(f"""
    __global__ void array_declaration(int *in_matrix, int *out_matrix, int in_integer)
    {{
      int d[{array_len}]; // 编译时已经是固定长度，符合语法
      out_matrix[0] = in_matrix[0];
    }}
    """)
  

• 方案三：如果每个线程需要独立的私有数组，直接在Host侧提前分配设备全局内存，作为参数传入kernel使用，虽然访存延迟高于栈和共享内存，但大小不受限制。

如果确实要尝试VLA写法，可以给SourceModule添加编译选项开启C99支持，但非常不推荐生产环境使用：

mod = SourceModule(""" 你的kernel代码 """, options=['--std=c99'])

内容的提问来源于stack exchange，提问作者Henry