关于CUDA浮点数运算结果不准确问题

最近在学习CudaC编程，在验证点积运算的过程中，发现GPU/CPU运算结果不一致问题。具体代码如下

#include "cuda_runtime.h"

#include "device_launch_parameters.h"

#include <time.h>

#include <iostream>

#include <iomanip>



#include <stdio.h>



const int cnTheadPerBlock = 256;

const int cnElementSz = 33*1024;



bool initialize(float* pfHostA, float* pfHostB, float& fHostC);



__global__ void dot(float* pfDeviceA, float* pfDeviceB, float* pfPartC) {

    __shared__ float fCache[cnTheadPerBlock];



    int nIDx = threadIdx.x + blockIdx.x * blockDim.x;

    int nCacheNo = threadIdx.x;



    float fTemp = 0.0;

    while (nIDx < cnElementSz) {

        fTemp += pfDeviceA[nIDx] * pfDeviceB[nIDx];

        nIDx += blockDim.x*gridDim.x;

    }

    fCache[nCacheNo] = fTemp;



    __syncthreads();



    int i = blockDim.x /2;

    while (i!=0) {

        if (nCacheNo < i)

            fCache[nCacheNo] += fCache[nCacheNo + i];

        __syncthreads();

        i /= 2;

    }

    if (nCacheNo == 0)

        pfPartC[blockIdx.x] = fCache[0];



}



int main()

{

    float afHostA[cnElementSz], afHostB[cnElementSz], fHostC;



    for (int n = 0; n < cnElementSz; n++) {

        afHostA[n] = static_cast<float>(n);

        afHostB[n] = static_cast<float>(2 * n);

    }



    clock_t tFormer1 = clock();



    initialize(afHostA, afHostB, fHostC);







    clock_t tPoster1 = clock();

    double dT1 = static_cast<double>(tPoster1 - tFormer1) / CLOCKS_PER_SEC;



    std::cout << "The result is:" << std::fixed << std::setprecision(0) << fHostC << std::endl;

    std::cout << "The captured time is:" << dT1  << std::endl;



    clock_t tFormer2 = clock();

    float fTempR = 0.0;

    for (int n = 0; n < cnElementSz; n++) {

        fTempR += 2.0f*static_cast<float>(n)* static_cast<float>(n);

    }

    clock_t tPoster2 = clock();

    double dT2 = static_cast<double>(tPoster2 - tFormer2) / CLOCKS_PER_SEC;



    std::cout << "The R is:"  << std::fixed << std::setprecision(0) << fTempR << std::endl;

    std::cout << "The capturedR time is:" << dT2 << std::endl;

    return 0;

}





bool initialize(float* pfHostA, float* pfHostB, float& fHostC) {

    float* pfDeviceA, * pfDeviceB, *pfPartC, *pfHostC;







    pfHostC = NULL;

    pfHostC = (float*)(malloc((cnElementSz / cnTheadPerBlock) * sizeof(float)));



    clock_t tMallocF = clock();

    if (cudaSuccess != cudaMalloc(((void**)&pfDeviceA), cnElementSz*sizeof(float)) ||

        cudaSuccess != cudaMalloc(((void**)&pfDeviceB), cnElementSz*sizeof(float)) ||

        cudaSuccess != cudaMalloc(((void**)&pfPartC), (cnElementSz/cnTheadPerBlock)*sizeof(float)) ) {

        std::cerr << "cudaMalloc failed!" << std::endl;

        return false;

    }

    clock_t tMallocP = clock();

    double dTimeML = static_cast<double>(tMallocP - tMallocF)/CLOCKS_PER_SEC;

    std::cout << "Malloc time is:" << dTimeML <<std::endl;



    clock_t tMemcpyF1 = clock();

    if (cudaSuccess != cudaMemcpy(pfDeviceA, pfHostA, cnElementSz * sizeof(float), cudaMemcpyHostToDevice) ||

        cudaSuccess != cudaMemcpy(pfDeviceB, pfHostB, cnElementSz * sizeof(float), cudaMemcpyHostToDevice)) {

        std::cerr << "cudaMemcpy failed!" << std::endl;

        return false;

    }

    clock_t tMemcpyP1 = clock();

    double dTimeMC1 = static_cast<double>(tMemcpyP1 - tMemcpyF1) / CLOCKS_PER_SEC;

    std::cout << "Memcpy1 time is:" << dTimeMC1 << std::endl;



    dot << <cnElementSz / cnTheadPerBlock, cnTheadPerBlock >> > (pfDeviceA, pfDeviceB, pfPartC);



    clock_t tMemcpyF2 = clock();

    if (cudaSuccess != cudaMemcpy(pfHostC, pfPartC, (cnElementSz / cnTheadPerBlock) * sizeof(float), cudaMemcpyDeviceToHost)) {

        std::cerr << "cudaMemcpy failed!" << std::endl;

        return false;

    }

    clock_t tMemcpyP2 = clock();

    double dTimeMC2 = static_cast<double>(tMemcpyP2- tMemcpyF2) / CLOCKS_PER_SEC;

    std::cout << "Memcpy2 time is:" << dTimeMC2 << std::endl;



    fHostC = 0.0;

    for (int n = 0; n < (cnElementSz / cnTheadPerBlock); n++) {

        fHostC += pfHostC[n];

    }



    free(pfHostC);

    cudaFree(pfDeviceA);

    cudaFree(pfDeviceB);

    cudaFree(pfPartC);



    return true;

}

在cnElemetSz小于256的情况下，CPU和GPU运算结果基本一致。但大于256的情况下CPU和GPU运算结果就不一样了。不知道是什么原因造成了这种情况，请大神不吝指教哈。下图是以上程序的执行结果。