visual-studio-2012 - カーネルから呼び出されたときに CUB DeviceScan を動作させる

Question

現在、CUB を使用してインテグレーター コードを書き直す方法を学ぼうとしています。ドキュメントの例とコード スニペットを見てきましたが、私がやろうとしていることの例をまだ見つけていません。具体的には、マスター スレッドから呼び出された InclusiveSum を実行します。私が見たものから、すべての例はデバイスではなくホストから関数を呼び出しますが、ここで実行できることを示唆しています:

「カーネル コードからこのメソッドを呼び出すときは、必ずコンパイラのマクロ定義で CUB_CDP マクロを定義してください。」

プロジェクトのプロパティ-> Cuda Linker->コマンドラインに移動し、「-DCUB_CDP」を追加して、Visual Studio 2012にこれを追加しようとしました。それが正しいかどうかはわかりませんが、次のビルド行を取得します。

"nvcc.exe" -gencode=arch=compute_35,code=\"sm_35,compute_35\" --use-local-env --cl-version 2012 -ccbin "C:\Program Files (x86)\Microsoft Visual Studio 11.0 \VC\bin\x86_amd64" -rdc=true -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v6.0\include" -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v6 .0\include" -G --keep-dir x64\Debug -maxrregcount=0 --machine 64 --compile -cudart static -DCUB_CDP -g -D_MBCS -Xcompiler "/EHsc /W3 /nologo /Od /Zi /RTC1 /MT " -o "x64\Debug\Algorithm Test.cu.obj" "C:\Users...\Algorithm Test.cu"

テスト コードには、実際のコードがどのように機能するかをシミュレートするために、1 つのスレッドで実行されるテスト カーネルが含まれます。

#define CUB_STDERR
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
#define NUMPOINTS 5*1024    
#define NUMTHREADSPERBLOCK 256
#define NUMBLOCKSPERGRID 32
#define MAXLENGTH NUMTHREADSPERBLOCK*NUMBLOCKSPERGRID   //Use multiple of 256

#include <cuda.h>
#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include <iostream>
#include <fstream>
#include <iomanip>                      //display 2 decimal places
#include <math.h>
#include <arrayFunctions.h>
#include <ctime>                        //For timers
#include <sstream>                      //For the filename
#include <assert.h>
#include <stdlib.h>
#include <cub/cub.cuh>


#if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ < 200)
#undef  assert
#define assert(arg)
#endif

__device__ __constant__ int numThreads = NUMTHREADSPERBLOCK;    //Number of threads per block
__device__ __constant__ int numBlocks = NUMBLOCKSPERGRID;       //Number of blocks per grid
__device__ __constant__ int maxlength = MAXLENGTH;  
__device__ double concSort[MAXLENGTH];

inline void gpuAssert(cudaError_t code, char *file, int line, bool abort=true)
{
    //Error checking
    if (code != cudaSuccess) 
    {
        fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
        if (abort) exit(code);
    }
}


using namespace std;
using namespace cub;

__global__ void test(double*);

int main(int argc, char** argv)
{
    cudaDeviceSetSharedMemConfig(cudaSharedMemBankSizeEightByte);
    cudaSetDevice(0);
    std::cout << std::fixed;                //Displays 2 decimal places.
    std::cout << std::setprecision(16);     //Displays 2 decimal places.

    const int maxlength = MAXLENGTH;        //Number of discrete concentrations tracking.
    double concs[maxlength] = {};           //Meant to store the initial concentrations .


    std::cout<<" ";
    std::cout<<"\n";    

    double *d_concs;                //The concentrations for a specific timestep.

    size_t size_concs = sizeof(concs);


    gpuErrchk(cudaMalloc((void**)&d_concs, size_concs));
    gpuErrchk(cudaMemcpy(d_concs, &concs, size_concs, cudaMemcpyHostToDevice));


    //Run the integrator.
    std::clock_t start;
    double duration;
    start = std::clock();

    test<<<1,1>>>(d_concs);

    std::cout<<"\n";

    gpuErrchk( cudaPeekAtLastError() );
    gpuErrchk( cudaDeviceSynchronize() );
    duration = (std::clock() - start)/ (double) CLOCKS_PER_SEC;
    std::cout<<"The calculation took this long: "<< duration <<'\n';

    std::cout<<"\n";

    gpuErrchk(cudaMemcpy(concs, d_concs, size_concs, cudaMemcpyDeviceToHost));
    cudaDeviceSynchronize();

    ///*
    for (int i=0; i < 33; i++)
    {
        std::cout << "\n";
        std::cout << concs[i];
    }
    //*/

    cudaDeviceReset();  //Clean up all memory.
    return 0;
}



__global__ void test(double* concs)
{
    int size=MAXLENGTH;
    int threads = NUMTHREADSPERBLOCK;
    int blocks = NUMBLOCKSPERGRID;

    for (int i = 0; i < size; i++)
        concs[i] = i * .00000000001;

    ///*
    void *d_temp_storage = NULL;
    size_t temp_storage_bytes = 0;
    CubDebug(cub::DeviceScan::InclusiveSum(d_temp_storage, temp_storage_bytes, concs, concs, size));
    cudaMalloc(&d_temp_storage, temp_storage_bytes);
    CubDebug(cub::DeviceScan::InclusiveSum(d_temp_storage, temp_storage_bytes, concs, concs, size));

}

次のエラーが表示されます。次の投稿から、このマクロ CUB_CDP の定義が私のエラーであることを示唆しています。

1>C:/Users/Karsten Chu/New Google Drive/Research/Visual Studio 2012/Projects/Dynamic Parallelism Test/Dynamic Parallelism Test/Algorithm Test.cu(146): error : calling a __host__ function("exit") from a __global__ function("test") is not allowed
1>C:/Users/Karsten Chu/New Google Drive/Research/Visual Studio 2012/Projects/Dynamic Parallelism Test/Dynamic Parallelism Test/Algorithm Test.cu(148): error : calling a __host__ function("exit") from a __global__ function("test") is not allowed

https://groups.google.com/forum/#!searchin/cub-users/CUB_CDP/cub-users/9ltP52Ohosg/uM9_RUy11e0J

このライブラリの使用方法を学ぶことは、物理学以外のことではなく、物理学に集中するのに本当に役立つと思うので、助けていただければ幸いです。