restrict 和模板函数默认参数的 nvcc 处理

Question

下面的代码是合法的 C++（用 g++ -Wall 编译干净）：

#ifdef MAKE_COMPILE
#define __restrict__ /* empty */
#define NO_FORWARD_DECLARATIONS
#endif

#include <stdio.h>

template <class T>
struct Array
{
int width, height;
T *ptr;
};

#ifdef HAVE_CUDA
template<typename T, int KernelSize>
     static __global__ void genConvolve_kernel(const T __restrict__ * inputImageArray , T __restrict__ * outputImageArray , int inputWidth , int outputWidth )
{
    if ((threadIdx.x == 4) && (threadIdx.y == 2))
       printf("Hello world from CUDA!\n");
}
#endif

#ifndef NO_FORWARD_DECLARATIONS
template <typename T, int KernelSize>
     void genConvolve_cuda(const Array<T> & kernelArray , const Array<T> & inputImageArray , Array<T> & outputImageArray , int blockWidth=16, int blockHeight=16);

template <typename T, int KernelSize>
     void genConvolve_cuda_deviceptrs( const T __restrict__ * inputImageArray , T __restrict__ * outputImageArray , int inputWidth , int outputWidth , int outputHeight , int blockWidth=16, int blockHeight=16);
#endif

template <typename T, int KernelSize>
     void genConvolve_cuda_deviceptrs( const T __restrict__ * inputImageArray , T __restrict__ * outputImageArray , int inputWidth , int outputWidth , int outputHeight , int blockWidth=16, int blockHeight=16)
{
#ifdef HAVE_CUDA
    dim3 block(blockWidth,blockHeight);
    dim3 grid(1,1);
    genConvolve_kernel<T,KernelSize><<<grid,block>>>(inputImageArray,outputImageArray,inputWidth,outputWidth);
#else
    printf("Hello, world!\n");
#endif
}

template <typename T, int KernelSize>
     void genConvolve_cuda( const Array<T> & kernelArray , const Array<T> & inputImageArray , Array<T> & outputImageArray , int blockWidth=16, int blockHeight=16)
{
    genConvolve_cuda_deviceptrs<T,KernelSize>((const T *)inputImageArray.ptr,outputImageArray.ptr, inputImageArray.width, outputImageArray.width, outputImageArray.height, blockWidth, blockHeight);
}

int main(int argc, char *argv[])
{
    Array<float> a;

    genConvolve_cuda<float,3>(a,a,a);
#ifdef HAVE_CUDA
    cudaDeviceSynchronize();
#endif

    return 0;
}

但是，当我尝试使用 nvcc 编译它时，出现错误：

nvcc t.cu

t.cu(39): warning: specifying a default argument when redeclaring an unreferenced function template is nonstandard

t.cu(39): warning: redefinition of default argument

t.cu(51): warning: specifying a default argument when redeclaring an unreferenced function template is nonstandard

t.cu(51): warning: redefinition of default argument

t.cu(53): error: template instantiation resulted in unexpected function type of "void (const float *, float *, int, int, int, int, int)" (the meaning of a name may have changed since the template declaration -- the type of the template is "void (const __restrict__ T *, __restrict__ T *, int, int, int, int, int)") detected during: instantiation of "genConvolve_cuda_deviceptrs" based on template arguments (53): here instantiation of "void genConvolve_cuda(const Array &, const Array &, Array &, int, int) [with T=float, KernelSize=3]" (60): here

（我在发布前清理示例时，行号略有偏移。）

当我定义 -DMAKE_COMPILE 时，警告和错误消失了；但是，我真的很想在头文件中指定前向声明，并使用 restrict !

所以两个问题：

当有默认函数参数（在我的例子中是 blockWidth 和 blockHeight？）时，如何使用 NVCC 指定模板函数的前向声明？
如何正确使用 __restrict__ 模板参数？

Answer 1

How to properly use __restrict__ with template arguments?

在与同事商量后，有人向我指出这个 __restrict__ 用法：

const T __restrict__ * inputImageArray ...

有疑问。为了使 __restrict__ 有任何效果，它应该放在星号和指针名称之间：

const T * __restrict__ inputImageArray ...

(gcc reference, and CUDA reference)

在您所展示的非标准用法中，gcc 似乎允许这样做但默默地 "drops" 意图； __restrict__ 的效果在这种情况下不适用。在这方面，CUDA 确实不同于 gcc 行为。但是因为如上所述是有问题的用法，所以不太可能将 nvcc 修改为 "fix" 这个问题。

如果您切换到标准 __restrict__ 用法，您可以使编译错误在您显示的代码中消失。如果您的意图是向编译器声明这些实际上是受限指针，无论如何建议这样做：

#ifdef MAKE_COMPILE
#define __restrict__ /* empty */
#define NO_FORWARD_DECLARATIONS
#endif

#include <stdio.h>

template <class T>
struct Array
{
int width, height;
T *ptr;
};

#ifdef HAVE_CUDA
template<typename T, int KernelSize>
     static __global__ void genConvolve_kernel(const T * __restrict__ inputImageArray , T * __restrict__ outputImageArray , int inputWidth , int outputWidth )
{
    if ((threadIdx.x == 4) && (threadIdx.y == 2))
       printf("Hello world from CUDA!\n");
}
#endif

#ifndef NO_FORWARD_DECLARATIONS
template <typename T, int KernelSize>
     void genConvolve_cuda(const Array<T> & kernelArray , const Array<T> & inputImageArray , Array<T> & outputImageArray , int blockWidth=16, int blockHeight=16);

template <typename T, int KernelSize>
     void genConvolve_cuda_deviceptrs( const T * __restrict__ inputImageArray , T * __restrict__ outputImageArray , int inputWidth , int outputWidth , int outputHeight , int blockWidth=16, int blockHeight=16);
#endif

template <typename T, int KernelSize>
     void genConvolve_cuda_deviceptrs( const T * __restrict__ inputImageArray , T * __restrict__ outputImageArray , int inputWidth , int outputWidth , int outputHeight , int blockWidth=16, int blockHeight=16)
{
#ifdef HAVE_CUDA
    dim3 block(blockWidth,blockHeight);
    dim3 grid(1,1);
    genConvolve_kernel<T,KernelSize><<<grid,block>>>(inputImageArray,outputImageArray,inputWidth,outputWidth);
#else
    printf("Hello, world!\n");
#endif
}

template <typename T, int KernelSize>
     void genConvolve_cuda( const Array<T> & kernelArray , const Array<T> & inputImageArray , Array<T> & outputImageArray , int blockWidth=16, int blockHeight=16)
{
    genConvolve_cuda_deviceptrs<T,KernelSize>((const T *)inputImageArray.ptr,outputImageArray.ptr, inputImageArray.width, outputImageArray.width, outputImageArray.height, blockWidth, blockHeight);
}

int main(int argc, char *argv[])
{
    Array<float> a;

    genConvolve_cuda<float,3>(a,a,a);
#ifdef HAVE_CUDA
    cudaDeviceSynchronize();
#endif

    return 0;
}

警告仍然存在；这似乎是一个单独的问题：

t986.cu(33)：警告：在重新声明未引用的函数模板时指定默认参数是非标准的

t986.cu(33)：警告：重新定义默认参数

t986.cu(45)：警告：在重新声明未引用的函数模板时指定默认参数是非标准的

t986.cu(45)：警告：重新定义默认参数

如果默认（模板）函数参数包含在第一个声明中而不是后续声明中，则可以使这些警告消失，如下所示：

#ifdef MAKE_COMPILE
#define __restrict__ /* empty */
#define NO_FORWARD_DECLARATIONS
#endif

#include <stdio.h>

template <class T>
struct Array
{
int width, height;
T *ptr;
};

#ifdef HAVE_CUDA
template<typename T, int KernelSize>
     static __global__ void genConvolve_kernel(const T * __restrict__ inputImageArray , T * __restrict__ outputImageArray , int inputWidth , int outputWidth )
{
    if ((threadIdx.x == 4) && (threadIdx.y == 2))
       printf("Hello world from CUDA!\n");
}
#endif

#ifndef NO_FORWARD_DECLARATIONS
template <typename T, int KernelSize>
     void genConvolve_cuda(const Array<T> & kernelArray , const Array<T> & inputImageArray , Array<T> & outputImageArray , int blockWidth=16, int blockHeight=16);

template <typename T, int KernelSize>
     void genConvolve_cuda_deviceptrs( const T * __restrict__ inputImageArray , T * __restrict__ outputImageArray , int inputWidth , int outputWidth , int outputHeight , int blockWidth=16, int blockHeight=16);
#endif

template <typename T, int KernelSize>
     void genConvolve_cuda_deviceptrs( const T * __restrict__ inputImageArray , T * __restrict__ outputImageArray , int inputWidth , int outputWidth , int outputHeight , int blockWidth, int blockHeight)
{
#ifdef HAVE_CUDA
    dim3 block(blockWidth,blockHeight);
    dim3 grid(1,1);
    genConvolve_kernel<T,KernelSize><<<grid,block>>>(inputImageArray,outputImageArray,inputWidth,outputWidth);
#else
    printf("Hello, world!\n");
#endif
}

template <typename T, int KernelSize>
     void genConvolve_cuda( const Array<T> & kernelArray , const Array<T> & inputImageArray , Array<T> & outputImageArray , int blockWidth, int blockHeight)
{
    genConvolve_cuda_deviceptrs<T,KernelSize>((const T *)inputImageArray.ptr,outputImageArray.ptr, inputImageArray.width, outputImageArray.width, outputImageArray.height, blockWidth, blockHeight);
}

int main(int argc, char *argv[])
{
    Array<float> a;

    genConvolve_cuda<float,3>(a,a,a);
#ifdef HAVE_CUDA
    cudaDeviceSynchronize();
#endif

    return 0;
}

虽然我同意这仍然不同于 g++ 行为。然而，gnu 工具在这里可能仍然是一个不寻常的例子。默认参数的重新定义仍然是意外的，而且 clang 和 cl.exe (microsoft) 都会有问题。

restrict 和模板函数默认参数的 nvcc 处理

nvcc handling of restrict and template function default arguments

templates

cuda

nvcc

__restrict__ 和模板函数默认参数的 nvcc 处理

nvcc handling of __restrict__ and template function default arguments

templates

cuda

nvcc

restrict 和模板函数默认参数的 nvcc 处理

nvcc handling of restrict and template function default arguments