0

GPUで次のことを行っています

float decPart = valAtIndex - (int)valAtIndex;
int docID = decPart * numDocs;

ここで、型の場合は valAtIdex 、型floatも numDocsfloatです。私の場合、decPart は 0.2 で、numDocs は 10 でした。ただし、docID を印刷すると、1 として印刷されます (2 のはずです)。誰かが私がどこで間違いを犯しているのか教えてもらえますか?

以下は、役立つ場合の完全な方法です

__global__ 
void finalNc(float* scSortedCounts, int* pos, int* maxCountEx, float numDocs, 
             int lengthStreamCompacted, int* finalNc, int actualLengthPos, 
             float* val, int* docIndex, int* acV ,int* ptwrite,int* diff,
             int* posIndex)
{ 
    int index = blockDim.x * blockIdx.x + threadIdx.x; 
    if(index < lengthStreamCompacted){ 
        float valAtIndex = scSortedCounts[index]; 
        float decPart = valAtIndex - (int)valAtIndex; 
        int docID = decPart * numDocs; 
        int actualCount = (int)valAtIndex; 
        int placeToWrite = maxCountEx[docID] + actualCount; 
        if( index == (lengthStreamCompacted -1 )){ 
            finalNc[placeToWrite] = actualLengthPos - pos[index]; 
        }else{ 
            finalNc[placeToWrite] = pos[index + 1] + pos[index]; 
        } 
    } 
}
4

1 に答える 1

3

あなたが言及した場合、あなたが見ている結果は正しいように見えます.混乱の原因は、計算の中間結果のIEEE単精度表現にあります(おそらく、それらの中間結果の印刷または表示で丸められます) ）。

あなたが提供する例では、値 0.2f はbinary32 値として正確に表現できません。2 つの可能な値は次のいずれかです。

3E4CCCCC (1.99999988079071044921875E-1)

また

3E4CCCCD (2.0000000298023223876953125E-1)

最初の値が の実際の値である場合、docID言及した中間計算は 1 を生成する必要があります (これが観察されたものです)。これが 2 番目の値の場合、結果は 2 になります。これは完全に予想される動作です。

IEEE 丸めモードの効果を説明し、ここでエラーがないことを安心させるために、次のコード例を見てください。これは、3 つの可能な変換のいずれかで求めている計算を実行しますfloat-integer単純な切り捨て、 IEEE 754 は負の無限大に向かって丸められ、IEEE 754 は正の無限大に向かって丸められます。カーネルをテンプレート化し、1 から 10 までのランダムな float 値のセットでケースごとに実行しました。自分でコンパイルして実行し、(コードのように) 単純な切り捨てが実際に意図したとおりに機能することを確認できます。出力での IEEE 準拠の「切り上げ」および「切り捨て」変換の動作と同様に。

#include <thrust/host_vector.h>
#include <thrust/device_ptr.h>
#include <thrust/device_malloc.h>
#include <thrust/generate.h>
#include <cstdlib>
#include <cstdio>

template<int version>
__global__
void kernel(const float *inputs, int *outputs, int numDocs, int N)
{
    int index = blockDim.x * blockIdx.x + threadIdx.x; 

    if(index < N){ 
        float valAtIndex = inputs[index];
        int intPart;
        switch(version) {
            case 2:
                intPart = __float2int_ru(valAtIndex);
                break;
            case 1:
                intPart = __float2int_rd(valAtIndex);
                break;
            case 0:
            default:
                intPart = int(valAtIndex);
                break;
        }
        float decPart = valAtIndex - intPart;
        int docID = decPart * numDocs;
        outputs[index] = docID;
    }
}

inline float frand(){
    return 1.0f + 9.0f * ((float)rand()/(float)RAND_MAX);
}

int main(void)
{
    const size_t N = 100;
    const int numdocs = 10;

    srand(time(NULL));

    thrust::host_vector<float> values(N);
    thrust::host_vector<int> outputs(3*N);
    std::generate(values.begin(), values.end(), frand);

    thrust::device_ptr<float> in = thrust::device_malloc<float>(N);
    float * _in = thrust::raw_pointer_cast(in);
    thrust::copy(values.begin(), values.end(), in);

    thrust::device_ptr<int> out = thrust::device_malloc<int>(3*N);
    int * _out = thrust::raw_pointer_cast(out);

    kernel<0><<<1,128>>>(_in, _out, numdocs, N);
    kernel<1><<<1,128>>>(_in, _out+N, numdocs, N);
    kernel<2><<<1,128>>>(_in, _out+(2*N), numdocs, N);

    thrust::copy(out, out+3*N, outputs.begin());

    for(int i=0; i<N; i++)
        printf("%.10f %d %d %d\n", 
                values[i], outputs[i], outputs[N+i], outputs[2*N+i]);

    return 0;
}
于 2012-06-30T09:22:04.123 に答える