对 cuda 性能有什么期望

Question

我最近使用 cuda 编写了一些简单的代码，但虽然我听说它可以加快速度，但我没有看到太多。 例如，我用 cuda 编写了 2 个代码（不使用任何库，如 NPP 和 ...）。

首先是一个模糊滤镜（没有纹理记忆）。

第二个基于模糊的分割（包括一些每个元素的乘法、加法、除法和一些求和，我使用了来自点积的示例代码和一些模糊过滤器）。

但在第一种情况下，我的代码有点慢！！！与类似的 OpenCV 代码相比，第二个代码只给了我 2 倍的速度。

我的系统（英特尔酷睿 i7 4700HQ 2.4GHz 和 Geforce GT 750m）（N46JV 华硕笔记本电脑） 我在 Windows 10 中使用 Visual Studio 2015、Cuda 8、OpenCV 3.1。

我的代码给了我结果并且逻辑上是正确的。

我的问题是我应该对我的系统有更多期望吗？

我也使用了一些 NPP 代码，但它并没有改变任何东西，它们甚至比我自己的代码慢一点（为了减少）

我之前问过一些人，这是我最后的希望！

例如求和码：

__global__ void GF3_kernel(double* oldU_temp,
int width,
int height,
double* d_partial_sum_t) {
__shared__ double L[nthreadsGF3];

//Initialize shared memory:
L[threadIdx.x] = 0;
__syncthreads();


//calculate pixel coordinates
const int tid_in = blockIdx.x * blockDim.x + threadIdx.x;
int tid_test = tid_in;

const int cacheXIndex = threadIdx.x;

//temperory sum
double temp = 0;


tid_test = tid_in;
//power random pixels of thread
while (tid_test < width * height) {
    temp += oldU_temp[tid_test];
    tid_test += gridDim.x * blockDim.x;
}
L[cacheXIndex] = temp;
__syncthreads();

int p_idx = blockDim.x / 2;
while (p_idx != 0) {
    if (cacheXIndex < p_idx) {
        L[cacheXIndex] += L[cacheXIndex + p_idx];
    }
    __syncthreads();
    p_idx /= 2;
}
if (cacheXIndex == 0) {
    d_partial_sum_t[blockIdx.x] = L[0];
}
}

内核调用将是：

    /////////////////////////////////////////////////////////// sum /////////////////////////////////////////////////////////////

    GF3_kernel << <gridGF3, nthreadsGF3 >> >(d_oldU_temp, width, height, d_partial_sum_t);


    // copy for final summation in cpu
    cudaDeviceSynchronize();
    cudaMemcpy(partial_sum_t, d_partial_sum_t, gridGF3 * sizeof(double),     cudaMemcpyDeviceToHost);
    cudaDeviceSynchronize();
    //Summation result
    sum_temp = 0;
    //calculate Summation
    for (int j = 0; j < gridGF3; j++) {
        sum_temp += partial_sum_t[j];
    }

类似的东西。

你怎么看？

图像是尺寸为 390x390 像素的 1 通道图像

Answer 1

您不能指望 GPGPU 计算能带来奇迹。如果你比较高端 cpus 和 gpus 的浮点峰值性能，你会得到大约 4 到 10 的因子，具体取决于精度。 如果你设法以某种有效的方式使用你的 CPU，你可以期望你可以实现这个加速，如果你可以充分利用 GPU 的潜力。这很难。首先，并不是所有的问题都适合大规模并行 gpu 架构，其次优化 gpu 代码很难。

如果您想尝试从 GPU 中榨取更多性能，请尝试优化您的内存访问。这就是关键，您将失去大部分性能。 全局内存中的非对齐访问和随机访问将花费您 90% 以上的性能。尽量对齐数据并将其缓存在共享内存中。这会给你带来很长的路要走。 另一个技巧是让内核使用更少的寄存器。一般来说，每个线程使用的寄存器越少，一个流式多处理器上可以同时调度的寄存器就越多。

__global__ void GF3_kernel(
double* oldU_temp,
int width,
int height,
double* d_partial_sum_t) {
__shared__ double L[nthreadsGF3];

//Initialize shared memory:
L[threadIdx.x] = 0;
__syncthreads(); //__syncthreads does hurt performance. Use it if absolutly needed


//calculate pixel coordinates

//const int tid_in = blockIdx.x * blockDim.x + threadIdx.x; //not needed
int tid_test = blockIdx.x * blockDim.x + threadIdx.x;

//const int cacheXIndex = threadIdx.x; // not needed

//temperory sum
double temp = 0;


//tid_test = tid_in; //unneeded assignment
//power random pixels of thread
const int endRange = width * height; //most likely done by compiler, but just in case its not.
const int step = gridDim.x * blockDim.x; //same here
while (tid_test < endRange) {
    temp += oldU_temp[tid_test]; // aligned access in global memory is fine
    tid_test += step;
}
L[cacheXIndex] = temp;
__syncthreads();


// This looks like an reduction to me
// Looks fine but there is some potential
// https://www.google.de/url?sa=t&rct=j&q=&esrc=s&source=web&cd=1&ved=0ahUKEwik5IO7nOzRAhWDDRoKHVUuC0kQFggcMAA&url=http%3A%2F%2Fdeveloper.download.nvidia.com%2Fcompute%2Fcuda%2F1.1-Beta%2Fx86_website%2Fprojects%2Freduction%2Fdoc%2Freduction.pdf&usg=AFQjCNFZmgdihbG17glvRmF-zPHfmUR4Aw&sig2=Yil-X_Byf6T6_dKAtN80Aw
int p_idx = blockDim.x / 2;
while (p_idx != 0) {
    if (cacheXIndex < p_idx) {
        L[cacheXIndex] += L[cacheXIndex + p_idx];
    }
    __syncthreads();
    p_idx /= 2;
}
if (cacheXIndex == 0) {
    d_partial_sum_t[blockIdx.x] = L[0];
}
}