2

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

这是内核

__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为所有 pthreads 调用声明，而 pthreads 调用的内核将获得最后一个的值cudaMemcpyToSymbol。是对的吗？

现在我试图在一个二维数组中改变我的常量内存。第二个维度将是之前的值，但第一个维度将是使用的 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) ) );

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

4

1 回答 1

2

是的，您应该能够以您想要的方式使用常量内存，但是cudaMemcpyToSymbol您使用的复制操作不正确。调用的第一个参数是一个符号，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 数组将按行主要顺序排列，因此您需要使用行间距乘以行索引作为每个复制操作的偏移量。

于 2013-01-31T08:11:14.887 回答