1

我正在尝试使用 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)]

有人可以帮我保持记忆吗

4

1 回答 1

4

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

__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];
    }
}

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

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

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

于 2013-10-07T06:11:38.293 回答