cuda 使用常量内存作为二维数组

Question

我在一个多线程“主机”程序中实现我的内核，其中每个主机线程都在调用内核。 我在使用常量内存时遇到了问题。在常量内存中会放置一些参数，但是对于每个线程它们是不同的。 我也在出现问题的地方构建了一个示例。

这是内核

__global__ void Kernel( int *aiOutput, int Length )
{
    int id = threadIdx.x + blockIdx.x * blockDim.x;

    int iValue = 0;

    // bound check
    if( id < Length )
    {
        if( id % 3 == 0 )
            iValue = c_iaCoeff[2];
        else if( id % 2 == 0 )
            iValue = c_iaCoeff[1];
        else
            iValue = c_iaCoeff[0];

        aiOutput[id] = iValue;
    }
    __syncthreads();
}

一个 pthread 正在调用这个函数。

void* WrapperCopy( void* params )
{
    // choose cuda device to perform on
    CUDA_CHECK_RETURN( cudaSetDevice( 0 ) );

    // cast of params
    SParams *_params = (SParams*)params;

    // copy coefficients to constant memory
    CUDA_CHECK_RETURN( cudaMemcpyToSymbol( c_iaCoeff, _params->h_piCoeff, 3*sizeof(int) ) );

    // loop kernel
    for( int i=0; i<100; i++ )
    {
        // perfrom kernel
        Kernel<<< BLOCKCOUNT, BLOCKSIZE >>>( _params->d_piArray, _params->iLength );
    }

    // copy data back from gpu
    CUDA_CHECK_RETURN( cudaMemcpy(
            _params->h_piArray, _params->d_piArray, BLOCKSIZE*BLOCKCOUNT*sizeof(int), cudaMemcpyDeviceToHost ) );

    return NULL;
}

常量内存就是这样声明的。

__constant__ int c_iaCoeff[ 3 ];

对于每个主机线程在h_piCoeff 中都有不同的值，并将其复制到常量内存中。

现在我对每个 pthread 调用都得到相同的结果，因为它们在c_iaCoeff 中都得到了相同的值。 我认为这是常量内存如何工作并且必须在上下文中声明的问题 - 在示例中，只有一个 c_iaCoeff 为所有 pthread 调用声明，而 pthread 调用的内核将获得最后一个 @ 的值987654327@。对吗？

现在我尝试在二维数组中更改我的常量内存。 第二个维度是之前的值，但第一个维度是使用的 pthread 的索引。

__constant__ int c_iaCoeff2[ THREADS ][ 3 ];

在内核中会以这种方式使用它。

iValue = c_iaCoeff2[iTId][2];

但我不知道是否可以以这种方式使用常量内存，是吗？ 当我尝试将数据复制到常量内存时也出现错误。

CUDA_CHECK_RETURN( cudaMemcpyToSymbol( c_iaCoeff[_params->iTId], _params->h_piCoeff, 3*sizeof(int) ) );

一般是否可以将常量内存用作二维数组，如果可以，我的失败在哪里？

Answer 1

是的，您应该能够以您想要的方式使用常量内存，但是您使用的cudaMemcpyToSymbol 复制操作不正确。调用的第一个参数是 symbol，API 会在运行时符号表中查找以获取您请求的常量内存符号的地址。因此不能将地址传递给调用（尽管您的代码实际上是将初始化的主机值传递给调用，但为什么我将把它作为练习留给读者）。

您可能错过了调用中可选的第四个参数，它是您请求的符号指向的内存的偏移量。因此，您应该能够执行以下操作：

cudaMemcpyToSymbol( c_iaCoeff,                    // symbol to lookup
                    _params->h_piCoeff,           // source location
                    3*sizeof(int),                // number of bytes to copy
                    (3*_params->iTId)*sizeof(int) // Offset in bytes
                   );

[标准免责声明：在浏览器中编写，未经测试。使用风险自负]

最后一个参数是从符号开始的字节偏移量。您的 2D 数组将按行主要顺序排列，因此您需要使用行间距乘以行索引作为每个复制操作的偏移量。