GT540M上CUDA代码性能低下

3
执行以下代码示例在GeForce GT540M上需要约750毫秒,而在GT330M上相同的代码执行时间为约250毫秒。
将dev_a和dev_b复制到CUDA设备内存需要在GT540M上花费约350毫秒,在GT330M上花费约250毫秒。执行“addCuda”并将其复制回主机在GT540M上需要另外约400毫秒,在GT330M上则为0毫秒。
这不是我预期的结果,所以我检查了设备的属性,并发现GT540M设备在除了多处理器数量外的每个方面都超过或等于GT330M - GT540M有2个,而GT330M有6个。这真的可能吗?如果是这样,它真的会对执行时间产生如此巨大的影响吗?
#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>
#include <stdlib.h>
#include <time.h>
#include <math.h>

#define T 512
#define N 60000*T

__global__ void addCuda(double *a, double *b, double *c) {
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    if(tid < N) {
        c[tid] = sqrt(fabs(a[tid] * b[tid] / 12.34567)) * cos(a[tid]);
    }
}

int main() {
    double *dev_a, *dev_b, *dev_c;

    double* a = (double*)malloc(N*sizeof(double));
    double* b = (double*)malloc(N*sizeof(double));
    double* c = (double*)malloc(N*sizeof(double));

    printf("Filling arrays (CPU)...\n\n");
    int i;
    for(i = 0; i < N; i++) {
        a[i] = (double)-i;
        b[i] = (double)i;
    }

    int timer = clock();
    cudaMalloc((void**) &dev_a, N*sizeof(double));
    cudaMalloc((void**) &dev_b, N*sizeof(double));
    cudaMalloc((void**) &dev_c, N*sizeof(double));
    cudaMemcpy(dev_a, a, N*sizeof(double), cudaMemcpyHostToDevice);
    cudaMemcpy(dev_b, b, N*sizeof(double), cudaMemcpyHostToDevice);

    printf("Memcpy time: %d\n", clock() - timer);
    addCuda<<<(N+T-1)/T,T>>>(dev_a, dev_b, dev_c);
    cudaMemcpy(c, dev_c, N*sizeof(double), cudaMemcpyDeviceToHost);

    printf("Time elapsed: %d\n", clock() - timer);

cudaFree(dev_a);
cudaFree(dev_b);
cudaFree(dev_c);
free(a);
free(b);
free(c);

return 0;
}

设备属性如下:
GT540M:
Major revision number:         2
Minor revision number:         1
Name:                          GeForce GT 540M
Total global memory:           1073741824
Total shared memory per block: 49152
Total registers per block:     32768
Warp size:                     32
Maximum memory pitch:          2147483647
Maximum threads per block:     1024
Maximum dimension 0 of block:  1024
Maximum dimension 1 of block:  1024
Maximum dimension 2 of block:  64
Maximum dimension 0 of grid:   65535
Maximum dimension 1 of grid:   65535
Maximum dimension 2 of grid:   65535
Clock rate:                    1344000
Total constant memory:         65536
Texture alignment:             512
Concurrent copy and execution: Yes
Number of multiprocessors:     2
Kernel execution timeout:      Yes

GT330M

Major revision number:         1
Minor revision number:         2
Name:                          GeForce GT 330M
Total global memory:           268435456
Total shared memory per block: 16384
Total registers per block:     16384
Warp size:                     32
Maximum memory pitch:          2147483647
Maximum threads per block:     512
Maximum dimension 0 of block:  512
Maximum dimension 1 of block:  512
Maximum dimension 2 of block:  64
Maximum dimension 0 of grid:   65535
Maximum dimension 1 of grid:   65535
Maximum dimension 2 of grid:   1
Clock rate:                    1100000
Total constant memory:         65536
Texture alignment:             256
Concurrent copy and execution: Yes
Number of multiprocessors:     6
Kernel execution timeout:      Yes

5
首先要说明的是,GT330M不支持双精度运算,因此您在一个设备上比较的是单精度结果,而在另一个设备上比较的是双精度结果。当前硬件上两者之间有8倍的性能差异。请问您的问题中还可以添加一下操作系统和CUDA版本吗?这些是移动设备,所以它们显然没有在同一主机上运行。 - talonmies
总的来说,我认为你需要在CUDA上进行更大规模的作业基准测试,以抵消设置/拆卸时间。虽然这张旧卡可以更有效地完成工作,但这真的是你实际应用中的因素吗? - Rup
机器1: GT540M(计算能力2.1) CUDA版本4.1 Intel Core i5-2410M Windows 7 64位。机器2: GT330M(计算能力1.2) CUDA版本4.1 Intel Core i5-520M 在Mac的bootcamp上运行Windows 7 64位。在CPU上计算类似的加法需要大约2500毫秒(在两台机器上都是如此)。我尝试将所有双精度替换为单精度,以查看是否有所改变,但结果并没有。即使应用程序不需要,GPU仍然使用双精度吗? - Thorkil Holm-Jacobsen
2个回答

2

我认为从设备到主机的复制不可能是 ~0ms(毫秒)。建议检查该复制是否存在问题。


主机到设备的复制出了些问题,我用来检查结果的函数也有问题。从规格上看,GT540M确实比预期更快。 - Thorkil Holm-Jacobsen

-3

看一下多处理器的数量。


1
很抱歉,这并没有解释任何事情。其中一款具有8个内核的6个MP每4个时钟周期只能为一个warp退役一个单精度FMAD。另一款具有48个内核的2个MP每2个时钟周期可以为两个warp退役一个单精度FMAD,加上它有20%更高的时钟频率。两种产品都可以执行有限的指令级并行处理,在GT300M的情况下只能进行潜在的单精度乘法,在GT540M中可以是单精度FMAD。 - talonmies

网页内容由stack overflow 提供, 点击上面的
可以查看英文原文,
原文链接