950Mよりも950M速いですか？

Question

私は現在、私の友人のためにいくつかのコードをテストしています。私はさまざまなMXMモジュールにアクセスしており、950Mと970Mの間でいくつかのコードをベンチマークする奇妙な結果を得ています。

int main(void) 
{ 
    static const int WORK_SIZE = 65530; 
    float *data = new float[WORK_SIZE]; 

    float time; 
    cudaEvent_t start, stop; 
    cudaEventCreate(&start); 
    cudaEventCreate(&stop); 

    for (int i = 0; i < 1000; i++) 
    { 
     cudaEventRecord(start, 0); 

     initialize (data, WORK_SIZE); 

     float *recCpu = cpuReciprocal(data, WORK_SIZE); 
     float *recGpu = gpuReciprocal(data, WORK_SIZE); 
     float cpuSum = std::accumulate (recCpu, recCpu+WORK_SIZE, 0.0); 
     float gpuSum = std::accumulate (recGpu, recGpu+WORK_SIZE, 0.0); 

     cudaEventRecord(stop, 0); 
     cudaEventSynchronize(stop); 
     cudaEventElapsedTime(&time, start, stop); 

     /* Verify the results */ 
     std::cout << i << ") gpuSum = "<<gpuSum<< " cpuSum = " <<cpuSum<< " time = " <<time<< std::endl; 

     delete[] recCpu; 
     delete[] recGpu; 
    } 

    /* Free memory */ 
    delete[] data; 

    return 0; 
} 

私は970Mは約9秒を平均化しながら、カードと950mの両方に3回は約7秒の平均テストを実行しました。さらに、CUDAツールキットに含まれているベンチマークプログラムを実行し、950で970対10GB/sの3GB /秒のH-D速度しか得られませんでした。これは、Ubuntu 14.04およびCUDA 7.5を実行しているSkylake Xeonで行われました。この矛盾について誰かに光を当てることはできますか？ cpuReciprocalとgpuReciprocalのコードは次のとおりです。

/** 
* CUDA kernel that computes reciprocal values for a given vector 
*/ 
__global__ void reciprocalKernel(float *data, unsigned vectorSize) { 
    unsigned idx = blockIdx.x*blockDim.x+threadIdx.x; 
    if (idx < vectorSize) 
     data[idx] = 1.0/data[idx]; 
} 

/** 
* Host function that copies the data and launches the work on GPU 
*/ 
float *gpuReciprocal(float *data, unsigned size) 
{ 
    float *rc = new float[size]; 
    float *gpuData; 

    CUDA_CHECK_RETURN(cudaMalloc((void **)&gpuData, sizeof(float)*size)); 
    CUDA_CHECK_RETURN(cudaMemcpy(gpuData, data, sizeof(float)*size, cudaMemcpyHostToDevice)); 

    static const int BLOCK_SIZE = 256; 
    const int blockCount = (size+BLOCK_SIZE-1)/BLOCK_SIZE; 
    reciprocalKernel<<<blockCount, BLOCK_SIZE>>> (gpuData, size); 

    CUDA_CHECK_RETURN(cudaMemcpy(rc, gpuData, sizeof(float)*size, cudaMemcpyDeviceToHost)); 
    CUDA_CHECK_RETURN(cudaFree(gpuData)); 
    return rc; 
} 

float *cpuReciprocal(float *data, unsigned size) 
{ 
    float *rc = new float[size]; 
    for (unsigned cnt = 0; cnt < size; ++cnt) rc[cnt] = 1.0/data[cnt]; 
    return rc; 
} 

Answer 1

ホストからデバイスへの帯域幅は、PCI Expressパフォーマンスによって与えられます。 950はPCI Express Gen 3 16xに接続されているようで、970は4xのようです。ハードウェアの設定が詳細ではない可能性が最も高い説明です。 nvidia-smiなどのNVIDIAツールは、この情報を提供します。 あなたのアルゴリズムに関しては、最も遅い部分はホストとの通信であり、次はデバイスメモリの帯域幅が約10倍高速で、最後は除算です。