表面存储器比全局存储器花费更多的时间（两倍）

Question

表面存储器比全局存储器花费更多的时间（两倍）

我正在优化cuda程序。因此，我首先从矩阵乘法程序的优化开始。我用于并行化的线程方案是Blocksize（1，1），Gridsize（N，N）。我将表面内存用于内存优化目的（因为此线程方案无法使用共享内存）。当我比较优化前后的时间时，我发现执行使用表面内存后需要花费两倍的时间（我尝试使用不同的线程方案，但问题仍然存在）。从我到目前为止所读的内容来看，全局内存比表面内存要慢。因此使用表面存储器应该花费更少的时间。下面我给出使用表面存储器的矩阵乘法程序。有人可以告诉我这是什么问题吗？

#include < stdio.h > 
#include < cuda.h >

//#define N 3

surface < void, 2 > a_surf;
surface < void, 2 > b_surf;
surface < void, 2 > c_surf;

void CUDA_SAFE_CALL(cudaError_t call, int line) {
    switch (call) {
    case cudaSuccess:
        break;
    default:
        printf("ERROR at line :%i.%d' ' %s\n",
            line, call, cudaGetErrorString(call));
        exit(-1);
        break;
    }

}

__global__ void mul(int N) {
    int a, b, c, temp;
    int i;

    unsigned int x = blockIdx.x * blockDim.x + (threadIdx.x);
    unsigned int y = blockIdx.y * blockDim.y + (threadIdx.y);
    if (x < N && y < N) {

        temp = 0;
        for (i = 0; i < N; i++) {
            surf2Dread( & a, a_surf, (x) * 4, i);
            surf2Dread( & b, b_surf, (i) * 4, y);
            temp += a * b;
        }
        c = temp;

        // Write to output surface
        surf2Dwrite(c, c_surf, x * 4, y);
    }
}

int main() {
    int N = 100;
    int a[N][N], b[N][N], c[N][N];
    int i, j;
    int temp;
    clock_t t1, t2;
    cudaArray * da, * db, * dc;
    cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc < int > ();

    dim3 dimBlock(1, 1);
    dim3 dimGrid(N, N);

    temp = 0;
    for (i = 0; i < N; i++)
        for (j = 0; j < N; j++)
            a[i][j] = ++temp;

    temp = 0;
    for (i = 0; i < N; i++)
        for (j = 0; j < N; j++)
            b[i][j] = ++temp;

    CUDA_SAFE_CALL(cudaMallocArray( & da, & channelDesc, N, N, cudaArraySurfaceLoadStore), __LINE__);
    CUDA_SAFE_CALL(cudaMallocArray( & db, & channelDesc, N, N, cudaArraySurfaceLoadStore), __LINE__);
    CUDA_SAFE_CALL(cudaMallocArray( & dc, & channelDesc, N, N, cudaArraySurfaceLoadStore), __LINE__);

    int s = N * N * sizeof(int);

    CUDA_SAFE_CALL(cudaMemcpyToArray(da, 0, 0, a, s, cudaMemcpyHostToDevice), __LINE__);
    CUDA_SAFE_CALL(cudaMemcpyToArray(db, 0, 0, b, s, cudaMemcpyHostToDevice), __LINE__);

    CUDA_SAFE_CALL(cudaBindSurfaceToArray(a_surf, da), __LINE__);
    CUDA_SAFE_CALL(cudaBindSurfaceToArray(b_surf, db), __LINE__);
    CUDA_SAFE_CALL(cudaBindSurfaceToArray(c_surf, dc), __LINE__);

    t1 = clock();
    mul <<<dimGrid, dimBlock>>> (N);
    t2 = clock();

    CUDA_SAFE_CALL(cudaMemcpyFromArray(c, dc, 0, 0, s, cudaMemcpyDeviceToHost), __LINE__);

    double t3 = (double) t2 - (double) t1;
    t3 = t3 / CLOCKS_PER_SEC;

    printf("\n CUDA time :%lf", t3);

    CUDA_SAFE_CALL(cudaFreeArray(da), __LINE__);
    CUDA_SAFE_CALL(cudaFreeArray(db), __LINE__);
    CUDA_SAFE_CALL(cudaFreeArray(dc), __LINE__);
}

Run Code Online (Sandbox Code Playgroud)

Answer 1

Rob*_*lla 5

优化缓存不是一件容易的事。因此，这种琐碎的概括如下：

从我到目前为止所读的内容来看，全局内存比表面内存要慢。因此，使用表面内存应该花费更少的时间。

在我看来，范围如此之广以至于是不正确的。这经常是正确的，但并非总是如此。细节很重要，适当的编程实践也很重要。

表面内存无非是带有中间缓存的全局内存。但是全局内存（在当前CUDA版本支持的所有GPU上）已经具有二级缓存（在某些情况下为一级缓存）的支持。

您提出的用于测试/比较的代码存在许多问题，我将指出：

您的计时方法不正确。这个：
```
t1 = clock();
mul <<<dimGrid, dimBlock>>> (N);
t2 = clock();
```
Run Code Online (Sandbox Code Playgroud)
将对内核启动的持续时间进行计时，而不是内核执行的持续时间。因此，这几乎永远不是正确的时间计时方法。我们可以通过cudaDeviceSynchronize();在计时区域中放置一个调用来解决此问题，以在计时关闭之前强制完成内核。
如果您对性能感兴趣，这是一个特别糟糕的构造：
```
dim3 dimBlock(1, 1);
```
Run Code Online (Sandbox Code Playgroud)
因为每个GPU扭曲中的每32个线程中有31个处于非活动状态，所以您未使用GPU的31/32的性能。这具有广泛的含义。我对研究这种情况的性能没有兴趣，并且您也不应该这样做（因为它不能反映编写良好的代码的实际性能），除非您对微基准测试（而不是比较基准测试）感兴趣。因此，您的代码应固定为每个块至少处理32个线程，理想情况下可处理256个或更多线程。
您没有提供“全局内存”比较用例。因此，我将提供一个。
您尚未说明许多其他对比较基准测试或性能分析重要的因素，例如您所运行的GPU和平台以及compile命令。
我认为问题规模太小。100x100矩阵的矩阵乘法位于代码的边缘，可以合理地占用GPU或测试其性能极限。因此，我将加大问题的规模。

关于问题大小参数，这对于缓存讨论很重要。首先，表面缓存倾向于是空间优化的缓存，而普通的L1和L2缓存是线性（缓存线）优化的。对于非常大的2D问题，表面缓存可能比L2具有更好的行为。但是对于很小的问题，差异将不那么明显。其次，表面高速缓存是L1和L2高速缓存的补充，因此一个好的优化策略是通过L1和L2 集中一些数据，并通过表面来集中其他数据，以最大化可用的高速缓存行。实际上，由于您的输入矩阵是只读的，因此进一步的优化可能是使用纹理而不是表面。但是从相反的角度来看，如果我的问题这么小为了完全适合L2缓存，则表面缓存不太可能带来重大改进。您最初的问题大小包括3个100x100 int数量的矩阵，因此每个矩阵约40 KB，或总计120K字节。这个问题的大小将适合大多数GPU的L2缓存。通过增加问题的大小（我们将看到-总计大约12MB），我们可以严重阻碍仅全局内存的情况。

这是一个经过充分处理的代码示例，已针对上述大多数问题进行了修改。当我在CUDA 7.5 / Fedora 20的Quadro5000 GPU上运行此代码时，我发现表面情况比全局内存情况快大约8倍：

$ cat t1129.cu
#include <stdio.h>
#include <iostream>

typedef int mytype;
const int blk_dim=16;

#define my_N 1000
#define A_VAL 1
#define B_VAL 2

surface < void, 2 > a_surf;
surface < void, 2 > b_surf;
surface < void, 2 > c_surf;

void CUDA_SAFE_CALL(cudaError_t call, int line) {
    switch (call) {
    case cudaSuccess:
        break;
    default:
        printf("ERROR at line :%i.%d' ' %s\n",
            line, call, cudaGetErrorString(call));
        exit(-1);
        break;
    }

}

#ifdef USE_GLOBAL
__global__ void mul(const mytype * __restrict__ d_a, const mytype * __restrict__ d_b, mytype * __restrict__ d_c, const int N)
#else
__global__ void mul(const int N)
#endif
{
    mytype a, b, c, temp;
    int i;

    unsigned int x = blockIdx.x * blockDim.x + (threadIdx.x);
    unsigned int y = blockIdx.y * blockDim.y + (threadIdx.y);
    if (x < N && y < N) {

        temp = 0;
        for (i = 0; i < N; i++) {
#ifdef USE_GLOBAL
            a = d_a[x*N+i];
            b = d_b[i*N+y];
#else
            surf2Dread( & a, a_surf, (x) * sizeof(mytype), i);
            surf2Dread( & b, b_surf, (i) * sizeof(mytype), y);
#endif
            temp += a * b;
        }
        c = temp;
#ifdef USE_GLOBAL
        d_c[x*N+y] = c;
#else
        // Write to output surface
        surf2Dwrite(c, c_surf, x * sizeof(mytype), y);
#endif
    }
}

int main() {
    const int N = my_N;
    mytype *a, *b, *c, *d_a, *d_b, *d_c;
    int i, j;
    clock_t t1, t2;
    cudaArray * da, * db, * dc;
    cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc < mytype > ();

    dim3 dimBlock(blk_dim, blk_dim);
    dim3 dimGrid((N+dimBlock.x-1)/dimBlock.x, (N+dimBlock.y-1)/dimBlock.y);
    int s = N * N * sizeof(mytype);

    a = (mytype *)malloc(s);
    b = (mytype *)malloc(s);
    c = (mytype *)malloc(s);

    CUDA_SAFE_CALL(cudaMalloc(&d_a, s), __LINE__);
    CUDA_SAFE_CALL(cudaMalloc(&d_b, s), __LINE__);
    CUDA_SAFE_CALL(cudaMalloc(&d_c, s), __LINE__);

    for (i = 0; i < N; i++)
        for (j = 0; j < N; j++)
            a[i*N+j] = A_VAL;

    for (i = 0; i < N; i++)
        for (j = 0; j < N; j++)
            b[i*N+j] = B_VAL;

    CUDA_SAFE_CALL(cudaMallocArray( & da, & channelDesc, N, N, cudaArraySurfaceLoadStore), __LINE__);
    CUDA_SAFE_CALL(cudaMallocArray( & db, & channelDesc, N, N, cudaArraySurfaceLoadStore), __LINE__);
    CUDA_SAFE_CALL(cudaMallocArray( & dc, & channelDesc, N, N, cudaArraySurfaceLoadStore), __LINE__);


    CUDA_SAFE_CALL(cudaMemcpyToArray(da, 0, 0, a, s, cudaMemcpyHostToDevice), __LINE__);
    CUDA_SAFE_CALL(cudaMemcpyToArray(db, 0, 0, b, s, cudaMemcpyHostToDevice), __LINE__);

    CUDA_SAFE_CALL(cudaBindSurfaceToArray(a_surf, da), __LINE__);
    CUDA_SAFE_CALL(cudaBindSurfaceToArray(b_surf, db), __LINE__);
    CUDA_SAFE_CALL(cudaBindSurfaceToArray(c_surf, dc), __LINE__);

#ifdef USE_GLOBAL
    CUDA_SAFE_CALL(cudaMemcpy(d_a, a, s, cudaMemcpyHostToDevice), __LINE__);
    CUDA_SAFE_CALL(cudaMemcpy(d_b, b, s, cudaMemcpyHostToDevice), __LINE__);
#endif
    t1 = clock();
#ifdef USE_GLOBAL
    mul <<<dimGrid, dimBlock>>> (d_a, d_b, d_c, N);
#else
    mul <<<dimGrid, dimBlock>>> (N);
#endif
    cudaDeviceSynchronize();
    t2 = clock();

    CUDA_SAFE_CALL(cudaMemcpyFromArray(c, dc, 0, 0, s, cudaMemcpyDeviceToHost), __LINE__);
#ifdef USE_GLOBAL
    CUDA_SAFE_CALL(cudaMemcpy(c, d_c, s, cudaMemcpyDeviceToHost), __LINE__);
#endif

    double t3 = (double) t2 - (double) t1;
    t3 = t3 / CLOCKS_PER_SEC;

    printf("\n CUDA time :%lf\n", t3);
    for (i=0; i < N*N; i++)
      if(c[i] != A_VAL*B_VAL*N) {std::cout << "mismatch at: " << i << ", was: " << c[i] << " should be: " << A_VAL*B_VAL*N << std::endl;  return 1;}

    CUDA_SAFE_CALL(cudaFreeArray(da), __LINE__);
    CUDA_SAFE_CALL(cudaFreeArray(db), __LINE__);
    CUDA_SAFE_CALL(cudaFreeArray(dc), __LINE__);
    std::cout << "Success!"  << std::endl;
    return 0;
}
[bob@cluster1 misc]$ nvcc -O3 -o t1129 t1129.cu
[bob@cluster1 misc]$ ./t1129

 CUDA time :0.028771
Success!
$ nvcc -O3 -DUSE_GLOBAL -o t1129 t1129.cu
$ ./t1129

 CUDA time :0.243635
Success!
$

Run Code Online (Sandbox Code Playgroud)

最后要说的是，还有许多其他我们可以讨论的优化，这些优化可能会以一种或另一种方式进行比较。但是，如果您实际上要执行快速矩阵乘法运算，则应使用CUBLAS。您不应该编写自己的矩阵乘法例程。

归档时间：	10 年，5 月前
查看次数：	243 次
最近记录：	10 年，5 月前