Tensorflow，GPU 与 CPU 相比快多少？ （余弦相似度计算）

Question

昨天我尝试使用 GPU 进行计算，但非常失望。我的想法是 GPU（拥有数千个核心）将比 CPU 至少快 1k。

我想先用GPU来计算余弦相似度，然后让我告诉你结果。

• 首先，GPU 无法计算成对相似度（矩阵到矩阵），它表示内存不足，尽管 GPU 的内存为 24 GB RAM（型号为 - 2x RTX 3090，88.8 tflops 24 GB ram）卡、64 GB RAM 系统）

所以我必须使用向量到矩阵计算，然后你需要进行迭代，例如，如果你的原始矩阵有 400k 行，那么这意味着你必须进行 400k 次迭代向量到矩阵的余弦相似度计算。

我使用了顺便说一句的张量流：

import tensorflow as tf
import numpy as np
Xn = np.random.uniform(0,10, (400000,100)).astype('float32')
X = tf.constant(Xn)

count = 1000 * 100
for i in range(count):
    tf.keras.losses.cosine_similarity(
        X[0],
        X,
        axis=-1
    )

如您所见，矩阵的形状为 400k，100 表示 400k 行 100 列。

然后我进行了100k以上的迭代。 这 100k 花费了 176 秒，是的，这比我笔记本上的 CPU 快了 20 倍，但说实话，你可以使用 python Numba 库实现相同的结果，该库将 python 代码转换为机器本机代码。 （但是 400k 迭代需要 30 分钟 - 使用 Numba，不知道使用 GPU 需要多长时间）

结果是，我很失望，我期望 GPU 在内存方面更强大，主要是至少比 CPU 快 1000 倍..我期望时间以毫或微秒的方式..

为什么他们仍然说这张卡具有 80 teraflops 等的能力，而你却用本机指令在 CPU 上同时计算它。主要是 - CPU 只有几个核心，GPU 有数千个核心..

我可以让它更快吗？

Answer 1

瓶颈在于每次迭代都要将大量数据复制到显卡或从显卡复制大量数据。当我在 OpenCL 中编写不进行任何复制的余弦相似度函数，并在 RTX4070 + RTX4060 ti 上分发简单的实现（总共~10000 个 CUDA 管道）时，它在68 秒内完成了 100k 次迭代。两张卡的 GPU 使用率均为 85%，并且它们的发热量比纯矩阵乘法低 10 摄氏度。代码如下，但仅用于基准测试，不适用于生产，未测试准确性，甚至未针对最佳性能进行优化（但可以通过查看维基百科页面来了解余弦相似度的定义进行简单尝试）：

（这里两张卡的理论最大值只有 ~50 TFLOPS）

#include <iostream>
#include <fstream>

#include "gpgpu.hpp"
int main()
{

    try
    {
        const int MAXVAL = 1024*16; 

        GPGPU::Computer computer(GPGPU::Computer::DEVICE_GPUS); 

        // cosine similarity for 400k vectors (each 100 elements) vs 1 vector
        computer.compile(
            
            R"(

            kernel void CosineSimilarity( 
                global float * dataPrediction,
                global float * similarity) 
            { 
                const int threadId=get_global_id(0); 
                const int localId = threadId % 100;                

                const int vectorId = threadId / 100;

                // cos sim = sum(AB) / (sqrt( sum(A^2) ) * sqrt( sum(B^2) ))
                
                // A = first vector
                local float A[100];
                // B = selected vector
                local float B[100];

                A[localId] = dataPrediction[localId];
                B[localId] = dataPrediction[vectorId * 100 + localId];
            

                // sum(AB)
                local float sumAB[100];
                // sum(A2)
                local float sumA2[100];
                // sum(B2)
                local float sumB2[100];

                sumAB[localId] = A[localId]*B[localId];
                sumA2[localId] = A[localId]*A[localId];
                sumB2[localId] = B[localId]*B[localId];
                barrier(CLK_LOCAL_MEM_FENCE);

                // totally unrolled reduction
                if(localId<50)
                {
                    sumAB[localId] += sumAB[localId + 50];
                    sumA2[localId] += sumA2[localId + 50];
                    sumB2[localId] += sumB2[localId + 50];
                }
                barrier(CLK_LOCAL_MEM_FENCE);
                if(localId<25)
                {
                    sumAB[localId] += sumAB[localId + 25];
                    sumA2[localId] += sumA2[localId + 25];
                    sumB2[localId] += sumB2[localId + 25];
                }
                barrier(CLK_LOCAL_MEM_FENCE);
                if(localId<12)
                {
                    sumAB[localId] += sumAB[localId + 12];
                    sumA2[localId] += sumA2[localId + 12];
                    sumB2[localId] += sumB2[localId + 12];
                }
                barrier(CLK_LOCAL_MEM_FENCE);
                if(localId==0)
                {
                    sumAB[0] += sumAB[24];
                    sumA2[0] += sumA2[24];
                    sumB2[0] += sumB2[24];
                }
                barrier(CLK_LOCAL_MEM_FENCE);
                if(localId<6)
                {
                    sumAB[localId] += sumAB[localId + 6];
                    sumA2[localId] += sumA2[localId + 6];
                    sumB2[localId] += sumB2[localId + 6];
                }
                barrier(CLK_LOCAL_MEM_FENCE);
                if(localId<3)
                {
                    sumAB[localId] += sumAB[localId + 3];
                    sumA2[localId] += sumA2[localId + 3];
                    sumB2[localId] += sumB2[localId + 3];
                }
                barrier(CLK_LOCAL_MEM_FENCE);
                if(localId==0)
                {
                    sumAB[0] += sumAB[1] + sumAB[2];
                    sumA2[0] += sumA2[1] + sumA2[2];
                    sumB2[0] += sumB2[1] + sumB2[2];
                }
                barrier(CLK_LOCAL_MEM_FENCE);

                const float sqrtSumA2 = sqrt(sumA2[0]);
                const float sqrtSumB2 = sqrt(sumB2[0]);
                const float result = sumAB[0] / (sqrtSumA2 * sqrtSumB2);

                // only 1 thread writes result (similarity of current vector)
                if(localId == 0)
                    similarity[vectorId*100] = result;

             })", "CosineSimilarity");
        

        // createArrayInput for copying from RAM to VRAM
        auto dataPrediction = computer.createArrayState<float>("dataPrediction",400000 * 100);

        // createArrayOutput for copying from VRAM to RAM
        auto similarity = computer.createArrayState<float>("similarity", 400000 * 100);

        auto kernelParams = dataPrediction.next(similarity);
        
        // benchmark for 20 times
        for (int i = 0; i < 20; i++)
        {
            size_t nanoSeconds;
            {
                GPGPU::Bench bench(&nanoSeconds);

                for(int j=0;j<100000;j++)
                    computer.compute(kernelParams, "CosineSimilarity", 0, 400000 * 100, 100);
                
                // 25th vector's similarity
                // float sim = similarity[25 * 100]
            }
            std::cout << nanoSeconds / 1000000000.0f << " seconds" << std::endl; 
        }
        

    }
    catch (std::exception& ex)
    {
        std::cout << ex.what() << std::endl; // any error is handled here
    }
    return 0;
}