C++ CLWrapper类代码示例

TEST(testcopybuffer, throwsifnotondevice) {
    if(!EasyCL::isOpenCLAvailable()) {
        cout << "opencl library not found" << endl;
        exit(-1);
    }
    cout << "found opencl library" << endl;

    EasyCL *cl = EasyCL::createForFirstGpuOtherwiseCpu();
    //CLKernel *kernel = cl->buildKernel("testeasycl.cl", "test");
    const int bufferSize = 100 * 1024 / 4;
    float *in = new float[bufferSize];
    float *in2 = new float[bufferSize];
    for(int i = 0; i < bufferSize; i++) {
        in[i] = i * 3;
        in2[i] = 23 + i;
    }
    CLWrapper *inwrapper = cl->wrap(bufferSize, in);
    CLWrapper *in2wrapper = cl->wrap(bufferSize, in2);
    inwrapper->copyToDevice();
//    in2wrapper->copyToDevice();

    bool threw = false;
    try {
        inwrapper->copyTo(in2wrapper);
    } catch(runtime_error &e) {
        threw = true;
    }
    EXPECT_TRUE(threw);
    delete cl;
}

void CopyBuffer::copy( EasyCL *cl, CLWrapper *sourceWrapper, int *target ) {
    // first we will copy it to another buffer, so we can copy it out
    int bufferSize = sourceWrapper->size();
//    float *copiedBuffer = new float[ bufferSize ];
    CLWrapper *targetWrapper = cl->wrap( bufferSize, target );
    targetWrapper->createOnDevice();

    // now copy it, via a kernel
    const string kernelSource = "\n"
        "kernel void copy( int N, global int const *source, global int *dest ) {\n"
          "  #define globalId ( get_global_id(0) )\n"
          "  if( (int)globalId < N ) {\n"
          "      dest[globalId] = source[globalId];\n"
          "  }\n"
       " }\n";
    CLKernel *kernel = cl->buildKernelFromString( kernelSource, "copy", "" );
    kernel->in( bufferSize )->in( sourceWrapper )->out( targetWrapper );
    int workgroupSize = 32;
    int numWorkgroups = ( bufferSize + workgroupSize - 1 ) / workgroupSize;
    kernel->run_1d( numWorkgroups * workgroupSize, workgroupSize );
    cl->finish();
    targetWrapper->copyToHost();

    delete targetWrapper;
    delete kernel;
//    delete[] copiedBuffer;
}

TEST( SLOW_testintwrapper_huge, testreadwrite ) {

    EasyCL *cl = EasyCL::createForFirstGpuOtherwiseCpu();
    CLKernel *kernel = cl->buildKernel("testeasycl.cl", "test_stress");
    const int N = 1000000;
    int *in = new int[N];
    for( int i = 0; i < N; i++ ) {
        in[i] = i * 3;
    }
    int *out = new int[N];
    CLWrapper *inwrapper = cl->wrap(N, in);
    CLWrapper *outwrapper = cl->wrap(N, out);
    inwrapper->copyToDevice();
    outwrapper->createOnDevice();
    kernel->input( inwrapper );
    kernel->output( outwrapper );
    int globalSize = N;
    int workgroupsize = cl->getMaxWorkgroupSize();
    globalSize = ( ( globalSize + workgroupsize - 1 ) / workgroupsize ) * workgroupsize;
    cout << "globalsize: " << globalSize << " workgroupsize " << workgroupsize << endl;
    kernel->run_1d( globalSize, workgroupsize );
    outwrapper->copyToHost();
    for( int i = 0; i < N; i++ ) {
       if( out[i] != 689514 ) {
           cout << "out[" << i << "] != 689514: " << out[i] << endl;
           exit(-1);
       }
    }

    delete outwrapper;
    delete inwrapper;
    delete kernel;
    delete cl;
}

THClStorage* THClStorage_newWithSize(THClState *state, int device, long size)
{
  THArgCheck(size >= 0, 2, "invalid size");

  if(size > 0)
  {
    StatefulTimer::timeCheck("THClStorage_newWithSize START");
    THClStorage *storage = (THClStorage*)THAlloc(sizeof(THClStorage));
    float *data = new float[size];
    storage->device = device;
    storage->cl = THClState_getClv2(state, storage->device);
    CLWrapper *wrapper = storage->cl->wrap( size, data );
    if(state->trace) cout << "new wrapper, size " << size << endl;
    if(state->trace) cout << "wrapper->createOnDevice()" << endl;
    wrapper->createOnDevice();
    storage->data = data;
    storage->wrapper = wrapper;

    storage->size = size;
    storage->refcount = 1;
    storage->flag = TH_STORAGE_REFCOUNTED | TH_STORAGE_RESIZABLE | TH_STORAGE_FREEMEM;
    StatefulTimer::timeCheck("THClStorage_newWithSize END");
    return storage;
  }
  else
  {
    return THClStorage_newv2(state, device);
  }
}

VIRTUAL void ActivationLayer::backward() {
    // have no weights to backprop to, just need to backprop the errors

//    CLWrapper *imagesWrapper = 0;
//    if( previousLayer->hasOutputWrapper() ) {
//        imagesWrapper = previousLayer->getOutputWrapper();
//    } else {
//        imagesWrapper = cl->wrap( previousLayer->getOutputSize(), previousLayer->getOutput() );
//        imagesWrapper->copyToDevice();
//    }

    CLWrapper *gradOutputWrapper = 0;
    bool weOwnGradOutputWrapper = false;
    if( nextLayer->providesGradInputWrapper() ) {
        gradOutputWrapper = nextLayer->getGradInputWrapper();
    } else {
        gradOutputWrapper = cl->wrap( getOutputSize(), nextLayer->getGradInput() );
        gradOutputWrapper->copyToDevice();
        weOwnGradOutputWrapper = true;
    }

    activationBackpropImpl->backward( batchSize, outputWrapper, gradOutputWrapper, gradInputWrapper );
//    gradInputCopiedToHost = false;

//    if( !previousLayer->hasOutputWrapper() ) {
//        delete imagesWrapper;
//    }
    if( weOwnGradOutputWrapper ) {
        delete gradOutputWrapper;
    }
}

TEST(testfloatwrapperconst, main) {
    if(!EasyCL::isOpenCLAvailable()) {
        cout << "opencl library not found" << endl;
        exit(-1);
    }
    cout << "found opencl library" << endl;

    EasyCL *cl = EasyCL::createForFirstGpuOtherwiseCpu();
    CLKernel *kernel = cl->buildKernelFromString(getKernel(), "test", "");
    float in[5];
    for(int i = 0; i < 5; i++) {
        in[i] = i * 3;
    }
    float out[5];
    CLWrapper *inwrapper = cl->wrap(5, (float const *)in);
    CLWrapper *outwrapper = cl->wrap(5, out);
    inwrapper->copyToDevice();
    kernel->input(inwrapper);
    kernel->output(outwrapper);
    kernel->run_1d(5, 5);
    outwrapper->copyToHost();
    assertEquals(out[0] , 7);
    assertEquals(out[1] , 10);
    assertEquals(out[2] , 13);
    assertEquals(out[3] , 16);
    assertEquals(out[4] , 19);
    cout << "tests completed ok" << endl;

    delete inwrapper;
    delete outwrapper;
    delete kernel;
    delete cl;
}

VIRTUAL void Adagrad::updateWeights(CLWrapper *weightsWrapper, CLWrapper *gradWeightsWrapper,
        AdagradState *trainerState) {

    int numWeights = trainerState->numWeights;
    float *working = new float[ numWeights ];
    CLWrapper *workingWrapper = cl->wrap(numWeights, working);
    workingWrapper->createOnDevice();

    CLMathWrapper clWeights(weightsWrapper);
    CLMathWrapper clGradWeights(gradWeightsWrapper);
    CLMathWrapper clSumSquares(trainerState->sumSquaresWrapper);
    CLMathWrapper clWorking(workingWrapper);

    // following all happens on gpu, via clmathwrapper:
    clWorking = clGradWeights;
    clWorking.squared();
    clSumSquares += clWorking;

    clWorking = clSumSquares;
    clWorking.sqrt();
    clWorking.inv();
    clWorking *= clGradWeights;
    clWorking *= - learningRate;
    clWeights += clWorking;

    delete workingWrapper;
    delete[] working;
}

VIRTUAL void PoolingLayer::forward() {
    CLWrapper *upstreamOutputWrapper = 0;
    if( previousLayer->hasOutputWrapper() ) {
        upstreamOutputWrapper = previousLayer->getOutputWrapper();
    } else {
        float *upstreamOutput = previousLayer->getOutput();
        upstreamOutputWrapper = cl->wrap( previousLayer->getOutputSize(), upstreamOutput );
        upstreamOutputWrapper->copyToDevice();
    }
    poolingForwardImpl->forward( batchSize, upstreamOutputWrapper, selectorsWrapper, outputWrapper );
    if( !previousLayer->hasOutputWrapper() ) {
        delete upstreamOutputWrapper;
    }

//    cout << "PoolingLayer::forward() selectors after forward: " << endl;
//    for( int i = 0; i < outputImageSize; i++ ) {
//        for( int j = 0; j < outputImageSize; j++ ) {
//            cout << selectors[ i * outputImageSize + j ] << " ";
//        }
//        cout << endl;
//    }

//    cout << "PoolingLayer::forward() selectorsWrapper after forward: " << endl;
//    PrintBuffer::printInts( cl, selectorsWrapper, outputImageSize, outputImageSize );
}

VIRTUAL void PoolingLayer::backward() {
    // have no weights to backprop to, just need to backprop the errors

    CLWrapper *gradOutputWrapper = 0;
    bool weOwnErrorsWrapper = false;
    if( nextLayer->providesGradInputWrapper() ) {
        gradOutputWrapper = nextLayer->getGradInputWrapper();
    } else {
        gradOutputWrapper = cl->wrap( getOutputSize(), nextLayer->getGradInput() );
        gradOutputWrapper->copyToDevice();
        weOwnErrorsWrapper = true;
    }

//    cout << "PoolingLayer::backward selectorsWrapper:" << endl;
//    PrintBuffer::printInts( cl, selectorsWrapper, outputImageSize, outputImageSize );

//    int *selectors = reinterpret_cast< int * >( selectorsWrapper->getHostArray() );
//    cout << "PoolingLayer::backward selectors before copy to host:" << endl;
//    for( int i = 0; i < outputImageSize; i++ ) {
//        for( int j = 0; j < outputImageSize; j++ ) {
//            cout << " " << selectors[i * outputImageSize + j];
//        }
//        cout << endl;
//    }
//    selectorsWrapper->copyToHost();
//    cout << "PoolingLayer::backward selectors after copy to host:" << endl;
//    for( int i = 0; i < outputImageSize; i++ ) {
//        for( int j = 0; j < outputImageSize; j++ ) {
//            cout << " " << selectors[i * outputImageSize + j];
//        }
//        cout << endl;
//    }
//    selectorsWrapper->copyToDevice();

//    selectorsWrapper->copyToHost();

    poolingBackpropImpl->backward( batchSize, gradOutputWrapper, selectorsWrapper, gradInputWrapper );

//    gradInputWrapper->copyToHost();
//    float *gradInput = reinterpret_cast< float * >( gradInputWrapper->getHostArray() );
//    cout << "gradInput:" << endl;
//    for( int i = 0; i < inputImageSize; i++ ) {
//        for( int j = 0; j < inputImageSize; j++ ) {
////            cout << " " << gradInput[i * inputImageSize + j];
//            if( gradInput[i * inputImageSize + j] != 0 ) {
//                cout << " *";
//            } else {
//                cout << " .";
//            }
//        }
//        cout << endl;
//    }

    if( weOwnErrorsWrapper ) {
        delete gradOutputWrapper;
    }
}

void propagateWithWipe( Propagate *prop, int batchSize, LayerDimensions dim, float *inputData, float *filters, float *biases, float *results ) {
    int inputDataSize = batchSize * dim.inputCubeSize;
    CLWrapper *dataWrapper = prop->cl->wrap( inputDataSize, inputData );
    dataWrapper->copyToDevice();

    int weightsSize = dim.filtersSize;
    CLWrapper *weightsWrapper = prop->cl->wrap( weightsSize, filters );
    weightsWrapper->copyToDevice();

    CLWrapper *biasWeightsWrapper = 0;
    if( dim.biased ) {
        biasWeightsWrapper = prop->cl->wrap( dim.numFilters, biases );
        biasWeightsWrapper->copyToDevice();
    }

    CLWrapper *resultsWrapper = prop->cl->wrap( batchSize * dim.outputCubeSize, results );
    memset( results, 99, sizeof(float) * batchSize * dim.outputCubeSize );
    resultsWrapper->copyToDevice(); // so we can wipe it...

    StatefulTimer::timeCheck("testpropagate: after data wrapper processing");
    prop->propagate( batchSize, dataWrapper, weightsWrapper, biasWeightsWrapper,
            resultsWrapper );
//    StatefulTimer::timeCheck("Propagate::propagate after call propagate");
    resultsWrapper->copyToHost();
//    StatefulTimer::timeCheck("Propagate::propagate after copytohost");
    delete resultsWrapper;

    delete dataWrapper;
    delete weightsWrapper;
    if( dim.biased ) {
        delete biasWeightsWrapper;
    }
}

void forwardWithWipe( Forward *prop, int batchSize, LayerDimensions dim, float *inputData, float *filters, float *biases, float *output ) {
    int inputDataSize = batchSize * dim.inputCubeSize;
    CLWrapper *dataWrapper = prop->cl->wrap( inputDataSize, inputData );
    dataWrapper->copyToDevice();

    int weightsSize = dim.filtersSize;
    CLWrapper *weightsWrapper = prop->cl->wrap( weightsSize, filters );
    weightsWrapper->copyToDevice();

    CLWrapper *biasWrapper = 0;
    if( dim.biased ) {
        biasWrapper = prop->cl->wrap( dim.numFilters, biases );
        biasWrapper->copyToDevice();
    }

    CLWrapper *outputWrapper = prop->cl->wrap( batchSize * dim.outputCubeSize, output );
    memset( output, 99, sizeof(float) * batchSize * dim.outputCubeSize );
    outputWrapper->copyToDevice(); // so we can wipe it...

    StatefulTimer::timeCheck("testforward: after data wrapper processing");
    prop->forward( batchSize, dataWrapper, weightsWrapper, biasWrapper,
            outputWrapper );
//    StatefulTimer::timeCheck("Forward::forward after call forward");
    outputWrapper->copyToHost();
//    StatefulTimer::timeCheck("Forward::forward after copytohost");
    delete outputWrapper;

    delete dataWrapper;
    delete weightsWrapper;
    if( dim.biased ) {
        delete biasWrapper;
    }
}

TEST(testcopybuffer, main) {
    if(!EasyCL::isOpenCLAvailable()) {
        cout << "opencl library not found" << endl;
        exit(-1);
    }
    cout << "found opencl library" << endl;

    EasyCL *cl = EasyCL::createForFirstGpuOtherwiseCpu();
    //CLKernel *kernel = cl->buildKernel("testeasycl.cl", "test");
    float in[5];
    float in2[5];
    for(int i = 0; i < 5; i++) {
        in[i] = i * 3.0f;
        in2[i] = 23.0f + i;
    }
    float out[5];
    CLWrapper *inwrapper = cl->wrap(5, in);
    CLWrapper *in2wrapper = cl->wrap(5, in2);
    CLWrapper *outwrapper = cl->wrap(5, out);
    inwrapper->copyToDevice();
    in2wrapper->copyToDevice();

    EXPECT_FALSE(in2wrapper->isDeviceDirty());
    inwrapper->copyTo(in2wrapper);
    EXPECT_TRUE(in2wrapper->isDeviceDirty());
  //  cl->finish();
    // check that in2 host-side unchanged:
    for(int i = 0; i < 5; i++) {
        in[i] = i * 3.0f;
        EXPECT_EQ(23.0f + i, in2[i]);
    }

    in2wrapper->copyToHost();
    // check that in2 is now a copy of in:
    for(int i = 0; i < 5; i++) {
        in[i] = i * 3.0f;
        EXPECT_EQ(i * 3.0f, in2[i]);
    }

    // check that modifying in2 doesnt modfiy in:
    in2[1] = 27;
    in2wrapper->copyToDevice();
    inwrapper->copyToHost();
    EXPECT_EQ(1 * 3.0f, in[1]);

    in2wrapper->copyToHost();
    EXPECT_EQ(1 * 3.0f, in[1]);
    EXPECT_EQ(27.0f, in2[1]);
    
    delete inwrapper;
    delete in2wrapper;
    delete outwrapper;
    delete cl;
}

VIRTUAL void ActivationPropagate::propagate( int batchSize, float *input, float *output ) {
//    cout << "ActivationPropagate::propagate( float * )" << endl;
    CLWrapper *inputWrapper = cl->wrap( getInputSize( batchSize ), input );
    CLWrapper *outputWrapper = cl->wrap( getResultsSize( batchSize ), output );

    inputWrapper->copyToDevice();
    propagate( batchSize, inputWrapper, outputWrapper );
    outputWrapper->copyToHost();    

    delete outputWrapper;
    delete inputWrapper;
}

VIRTUAL void ActivationLayer::forward() {
    CLWrapper *inputWrapper = 0;
    if( previousLayer->hasOutputWrapper() ) {
        inputWrapper = previousLayer->getOutputWrapper();
    } else {
        float *input = previousLayer->getOutput();
        inputWrapper = cl->wrap( previousLayer->getOutputSize(), input );
        inputWrapper->copyToDevice();
    }
    activationForwardImpl->forward( batchSize, inputWrapper, outputWrapper );
//    outputCopiedToHost = false;
    if( !previousLayer->hasOutputWrapper() ) {
        delete inputWrapper;
    }
}

TEST( testMemset, basic ) {
    EasyCL *cl = DeepCLGtestGlobals_createEasyCL();

    CLKernel *kMemset = 0;
    // [[[cog
    // import stringify
    // stringify.write_kernel2( "kMemset", "cl/memset.cl", "cl_memset", '""' )
    // ]]]
    // generated using cog, from cl/memset.cl:
    const char * kMemsetSource =  
    "// Copyright Hugh Perkins 2015 hughperkins at gmail\n"
    "//\n"
    "// This Source Code Form is subject to the terms of the Mozilla Public License,\n"
    "// v. 2.0. If a copy of the MPL was not distributed with this file, You can\n"
    "// obtain one at http://mozilla.org/MPL/2.0/.\n"
    "\n"
    "kernel void cl_memset(global float *target, const float value, const int N) {\n"
    "    #define globalId get_global_id(0)\n"
    "    if ((int)globalId < N) {\n"
    "        target[globalId] = value;\n"
    "    }\n"
    "}\n"
    "\n"
    "";
    kMemset = cl->buildKernelFromString(kMemsetSource, "cl_memset", "", "cl/memset.cl");
    // [[[end]]]

    int N = 10000;
    float *myArray = new float[N];
    CLWrapper *myArrayWrapper = cl->wrap( N, myArray );
    myArrayWrapper->createOnDevice();
    kMemset->out( myArrayWrapper )->in( 99.0f )->in( N );
    int workgroupSize = 64;
    kMemset->run_1d( ( N + workgroupSize - 1 ) / workgroupSize * workgroupSize, workgroupSize );
    cl->finish();
    myArrayWrapper->copyToHost();
    for( int i = 0; i < 10; i++ ) {
//        cout << "myArray[" << i << "]=" << myArray[i] << endl;
    }
    for( int i = 0; i < N; i++ ) {
        EXPECT_EQ( 99.0f, myArray[i] );
    }

    delete kMemset;

    delete cl;
}