Получение CUB DeviceScan для работы при вызове из ядра

В настоящее время я пытаюсь узнать, как использовать CUB, чтобы, возможно, переписать код моего интегратора. Я просматривал примеры и фрагменты кода в документах, но я еще не нашел пример того, что я пытаюсь сделать. В частности, это для запуска InclusiveSum, вызываемого из основного потока. Из того, что я видел, все примеры вызывают функцию с хоста, а не устройства, но намекают, что это можно сделать здесь: http://nvlabs.github.io/cub/structcub_1_1_device_scan.html

"При вызове этого метода из кода ядра обязательно определите макрос CUB_CDP в определениях макросов вашего компилятора".

Я попытался добавить это в Visual Studio 2012, перейдя в свойства моего проекта->Cuda Linker-> Командная строка и добавив "-DCUB_CDP". Я не уверен, что это правильно, но я получаю следующую строку сборки:

"nvcc.exe" -gencode = arch = compute_35, code = \ "sm_35, compute_35 \" --use-local-env --cl-версия 2012 -ccbin "C:\Program Files (x86)\Microsoft Visual Studio 11.0\VC\bin\x86_amd64" -rdc=true -I"C:\Program Files\ Компьютерный инструментарий NVIDIA GPU \CUDA\v6.0\include" -I"C:\Program Files\ Компьютерный инструментарий NVIDIA GPU \CUDA\v6.0\include" -G --keep-dir x64\Debug -maxrregcount=0 - машина 64 - компилировать -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

Я был бы признателен за любую помощь, так как я думаю, что изучение того, как использовать эту библиотеку, могло бы реально помочь мне сосредоточиться на ФИЗИКЕ, а не на чем-либо, кроме физики.

Источник

1 ответ

Решение

Удалить CubDebugExit() Оболочка от вашего детёныша вызывает в тестовом ядре. Тогда ваш код скомпилируется.

Вместо этого:

CubDebugExit(cub::DeviceScan::ExclusiveSum(d_temp_storage, temp_storage_bytes, concs, concSort, maxlength));
cudaMalloc(&d_temp_storage, temp_storage_bytes);
CubDebugExit(cub::DeviceScan::ExclusiveSum(d_temp_storage, temp_storage_bytes, concs, concSort, maxlength));

Сделай это:

cub::DeviceScan::ExclusiveSum(d_temp_storage, temp_storage_bytes, concs, concSort, maxlength);
cudaMalloc(&d_temp_storage, temp_storage_bytes);
cub::DeviceScan::ExclusiveSum(d_temp_storage, temp_storage_bytes, concs, concSort, maxlength);

Макрос CubDebugExit не может использоваться в коде устройства.

Если вы предпочитаете, вы также можете использовать CubDebug() вместо CubDebugExit() обертка / макро.

Источник
Другие вопросы по тегам