c++ - CUDA 常量内存错误

Question

我正在尝试使用 CUDA 5.5 编写具有常量内存的示例代码。我有 2 个大小为 3000 的常量数组。我有另一个大小为 N 的全局数组 X。我要计算

Y[tid] = X[tid]*A[tid%3000] + B[tid%3000]

这是代码。

#include <iostream>
#include <stdio.h>
using namespace std;

#include <cuda.h>



__device__ __constant__ int A[3000];
__device__ __constant__ int B[3000];


__global__ void kernel( int *dc_A, int *dc_B, int *X, int *out, int N)
{
    int tid = threadIdx.x + blockIdx.x*blockDim.x;
    if( tid<N )
    {
        out[tid] = dc_A[tid%3000]*X[tid] + dc_B[tid%3000];
    }

}

int main()
{
    int N=100000;

    // set affine constants on host
    int *h_A, *h_B ; //host vectors
    h_A = (int*) malloc( 3000*sizeof(int) );
    h_B = (int*) malloc( 3000*sizeof(int) );
    for( int i=0 ; i<3000 ; i++ )
    {
        h_A[i] = (int) (drand48() * 10);
        h_B[i] = (int) (drand48() * 10);
    }

    //set X and Y on host
    int * h_X = (int*) malloc( N*sizeof(int) );
    int * h_out = (int *) malloc( N*sizeof(int) );
    //set the vector
    for( int i=0 ; i<N ; i++ )
    {
        h_X[i] = i;
        h_out[i] = 0;
    }

    // copy, A,B,X,Y to device
    int * d_X, *d_out;
    cudaMemcpyToSymbol( A, h_A, 3000 * sizeof(int) ) ;
    cudaMemcpyToSymbol( B, h_B, 3000 * sizeof(int) ) ;

    cudaMalloc( (void**)&d_X, N*sizeof(int) ) );
    cudaMemcpy( d_X, h_X, N*sizeof(int), cudaMemcpyHostToDevice ) ;
    cudaMalloc( (void**)&d_out, N*sizeof(int) ) ;



    //call kernel for vector addition
    kernel<<< (N+1024)/1024,1024 >>>(A,B, d_X, d_out, N);
    cudaPeekAtLastError() ;
    cudaDeviceSynchronize() ;


    // D --> H
    cudaMemcpy(h_out, d_out, N * sizeof(int), cudaMemcpyDeviceToHost ) ;


    free(h_A);
    free(h_B);


    return 0;
}

我正在尝试对这段代码运行调试器进行分析。事实证明，在复制到常量内存的行上，调试器出现以下错误

Coalescing of the CUDA commands output is off.
[Thread debugging using libthread_db enabled]
[New Thread 0x7ffff5c5b700 (LWP 31200)]

谁能帮我解决内存问题

Answer 1

这里有几个问题。从展示使用这两个常量数组的“正确”方法开始可能更容易，然后解释为什么你所做的不起作用。所以内核应该是这样的:

__global__ void kernel(int *X, int *out, int N)
{
    int tid = threadIdx.x + blockIdx.x*blockDim.x;
    if( tid<N )
    {
        out[tid] = A[tid%3000]*X[tid] + B[tid%3000];
    }
}

即。不要尝试将 A 和 B 传递给内核。原因如下:

1. 有些令人困惑的是，主机代码中的A 和B 不是有效的设备内存地址。它们是主机符号，提供运行时设备符号查找的 Hook 。将它们传递给内核是非法的 - 如果您想要它们的设备内存地址，则必须使用 cudaGetSymbolAddress 在运行时检索它。
2. 即使您确实调用了 cudaGetSymbolAddress 并在常量内存中检索符号设备地址，您也不应将它们作为参数传递给内核，因为这样做不会在运行内核。正确使用常量内存需要编译器发出特殊的 PTX 指令，并且编译器只会在知道特定的全局内存位置在常量内存中时才会这样做。如果您按值传递一个常量内存地址作为参数，则 __constant__ 属性会丢失，编译器无法知道生成正确的加载指令

一旦你开始工作，你会发现它非常慢，如果你分析它，你会发现指令重放和序列化程度非常高。使用常量内存的整个想法是，当 warp 中的每个线程都访问常量内存中的相同值 时，您可以利用常量缓存广播机制。您的示例与此完全相反 - 每个线程都在访问不同的值。在这种用例中，常规全局内存会更快。另请注意，模运算符在当前 GPU 上的性能很差，您应尽可能避免使用它。