diff --git a/src/hurricanedata/datareader.cu b/src/hurricanedata/datareader.cu index 4534aa4..f62f316 100644 --- a/src/hurricanedata/datareader.cu +++ b/src/hurricanedata/datareader.cu @@ -44,6 +44,9 @@ void DataReader::loadFile(T* dataOut, size_t fileIndex) { template void DataReader::loadFile(T* dataOut, size_t fileIndex, const string& varName) { + // TODO: We could make the index wrap around so that we dont lose data + // using fileIndex % NUMBEROFFILES here. + NcFile data(filePathManager.getPath(fileIndex), NcFile::read); multimap vars = data.getVars(); diff --git a/src/hurricanedata/datareader.h b/src/hurricanedata/datareader.h index 209132f..9c82aea 100644 --- a/src/hurricanedata/datareader.h +++ b/src/hurricanedata/datareader.h @@ -8,17 +8,42 @@ #include "filepathmanager.h" +/** + * @brief Simple wrapper around all the NetCDF functionality. + */ class DataReader { public: + /** + * @brief Constructor. + * @param path Path to the directory containing all the .nc4 files. + * @param variableName The variable we are interested in. + */ DataReader(const std::string &path, std::string variableName); + + /** + * @brief The length of the flat data in variableName. + * Used for allocating the right amount of memory. + */ size_t fileLength(size_t fileIndex); + /** + * @brief Write all the data in file fileIndex of variable we re interested in into the dataOut. + */ template void loadFile(T* dataOut, size_t fileIndex); + /** + * @brief Write all the data in file fileIndex of variable variableName into the dataOut. + * @param dataOut pointer to memory that should be written to. + * @param fileIndex the index of the file we want to load into memory. + * @param variableName the name of the variable + */ template void loadFile(T* dataOut, size_t fileIndex, const std::string& variableName); + /** + * @brief Get size of a variable. + */ size_t axisLength(size_t fileIndex, const std::string& axisName); ~DataReader(); diff --git a/src/hurricanedata/fielddata.h b/src/hurricanedata/fielddata.h index d7ac532..4a3257c 100644 --- a/src/hurricanedata/fielddata.h +++ b/src/hurricanedata/fielddata.h @@ -7,6 +7,7 @@ struct FieldMetadata { size_t widthSize; // Number of different longitudes size_t heightSize; // Number of different latitudes size_t depthSize; // Number of different levels + size_t timeSize; // Number of different times // lons is a managed Unified Memory array of size widthCount that indicates // that getVal(t, i, j, k) is a value with longitude of lons[i]. @@ -14,26 +15,37 @@ struct FieldMetadata { double *lons; double *lats; double *levs; - - size_t timeSize; // Number of different times - - // times is a managed Unified Memory array of size numberOfTimeStepsPerFile int *times; + size_t numberOfTimeStepsPerFile; }; using FieldMetadata = FieldMetadata; +/** + * @brief Allows for accessing data of a time-slice of a scalar field + * that may start in the middle of a file or go range over multiple files + * by holding references to multiple files at a time. + * + * @note Use the getVal method to index into it and get values. + */ struct FieldData { static constexpr size_t FILESNUM = 2; // Number of files stored in a FieldData struct. + static constexpr size_t numberOfTimeStepsPerField = 2; + + size_t fieldInd; // Indicates // A uniform array of length FILESNUM storing pointers to 4D arrays stored in device memory. float **valArrays; - size_t fieldInd; }; - +/** + * @brief Get the scalar field value at a particular index (timeInd, levInd, latInd, lonInd). + * Note that the timeInd may be counter-intuitive. See explanation in gpubufferhandler.h. + * + * @return scalar field float value. + */ extern __device__ float getVal( const FieldMetadata &md, const FieldData &d, diff --git a/src/hurricanedata/filepathmanager.h b/src/hurricanedata/filepathmanager.h index 2b1304d..96f9ec5 100644 --- a/src/hurricanedata/filepathmanager.h +++ b/src/hurricanedata/filepathmanager.h @@ -4,6 +4,15 @@ #include #include +/** + * @brief Simple class that is responsible for mapping a file index to a file path. + * So for example: + * index = 0 -> MERRA2_400.inst6_3d_ana_Np.20120101.nc4 + * index = 1 -> MERRA2_400.inst6_3d_ana_Np.20120102.nc4 + * index = 2 -> MERRA2_400.inst6_3d_ana_Np.20120103.nc4 + * index = 3 -> MERRA2_400.inst6_3d_ana_Np.20120104.nc4 + * etc... + */ class FilePathManager { public: FilePathManager(const std::string& path); diff --git a/src/hurricanedata/gpubuffer.h b/src/hurricanedata/gpubuffer.h index aaf1768..b953abe 100644 --- a/src/hurricanedata/gpubuffer.h +++ b/src/hurricanedata/gpubuffer.h @@ -10,18 +10,44 @@ struct DataHandle { float *d_data; // Device memory size_t size; + // Could include the data stored in host memory h_data in this handle if it were needed. }; +/** + * @brief Handles the asynchronous (un)loading data to the GPU. The rest of the + * application should not have to interface directly with this class. Getting data + * should go over the GPUBufferHandler class. + * + * @note This class uses a queue to give loading jobs to the worker thread. + * NetCDF-C is not thread safe so you may never read data using netCDF directly + * on any other thread than this worker thread. + * + * Assumes all the data in the nc4 files is the same size. + * + */ class GPUBuffer { public: - static constexpr size_t numBufferedFiles = 3; + static constexpr size_t numBufferedFiles = 3; // Number of files stored in memory at one time. GPUBuffer(DataReader& dataReader); + /** + * @brief Asynchronously tells the GPUBuffer to eventually load a particular file index + * into a buffer index (in which part of the buffer the file should be stored). + */ void loadFile(size_t fileIndex, size_t bufferIndex); // No blocking + /** + * @brief Get the values stored in a particular buffer index + * @return A struct DataHandle that points to the memory and gives its size. + */ DataHandle getDataHandle(size_t bufferIndex); // Potentially blocking + /** + * @brief Get the data from an axis (e.g. longitude) and its size. + * @note This is a blocking operation, so it makes a job for the worker + * to read the data and then waits untill the job is completed. + */ template std::pair getAxis(size_t fileIndex, const std::string& axisName); // Most probably blocking diff --git a/src/hurricanedata/gpubufferhandler.cu b/src/hurricanedata/gpubufferhandler.cu index eca628e..606a47c 100644 --- a/src/hurricanedata/gpubufferhandler.cu +++ b/src/hurricanedata/gpubufferhandler.cu @@ -26,7 +26,7 @@ gpuBuffer(gpuBuffer), fieldInd(0), bufferInd(0), fileInd(0) { fileInd++; } - fmd->timeSize = GPUBufferHandler::numberOfTimeStepsPerField; + fmd->timeSize = FieldData::numberOfTimeStepsPerField; cudaMallocManaged(&fmd->times, sizeof(fmd->numberOfTimeStepsPerFile*sizeof(int))); @@ -58,8 +58,8 @@ FieldData GPUBufferHandler::nextFieldData() { size_t newFieldInd = (fieldInd + 1) % fmd->numberOfTimeStepsPerFile; size_t newBufferInd = (bufferInd + ((fieldInd + 1) / fmd->numberOfTimeStepsPerFile)) % GPUBuffer::numBufferedFiles; - size_t endFieldInd = (fieldInd + GPUBufferHandler::numberOfTimeStepsPerField - 1) % fmd->numberOfTimeStepsPerFile; - size_t endBufferInd = (bufferInd + (fieldInd + GPUBufferHandler::numberOfTimeStepsPerField - 1)/fmd->numberOfTimeStepsPerFile) % GPUBuffer::numBufferedFiles; + size_t endFieldInd = (fieldInd + FieldData::numberOfTimeStepsPerField - 1) % fmd->numberOfTimeStepsPerFile; + size_t endBufferInd = (bufferInd + (fieldInd + FieldData::numberOfTimeStepsPerField - 1)/fmd->numberOfTimeStepsPerFile) % GPUBuffer::numBufferedFiles; size_t newEndFieldInd = (endFieldInd + 1) % fmd->numberOfTimeStepsPerFile; size_t newEndBufferInd = (endBufferInd + ((endFieldInd + 1) / fmd->numberOfTimeStepsPerFile)) % GPUBuffer::numBufferedFiles; diff --git a/src/hurricanedata/gpubufferhandler.h b/src/hurricanedata/gpubufferhandler.h index 2f1cfcc..2b4338b 100644 --- a/src/hurricanedata/gpubufferhandler.h +++ b/src/hurricanedata/gpubufferhandler.h @@ -6,18 +6,48 @@ #include +/** + * @brief Responsible for deciding when the GPUBuffer should load and unload files. + * Also assembles and gives access to FieldData. + * + * You will need to interface with this class. + */ class GPUBufferHandler { public: GPUBufferHandler(GPUBuffer& gpuBuffer); + /** + * @brief Produces a FieldData which can be used to retrieve values for a time-slice + * into a scalar field with a width of FieldData::numberOfTimeStepsPerField. + * + * @details This method always increments the start point of the time-slice + * by 1. See below: + * + * time steps = 0 1 2 3 4 5 6 7 8 9 10 11 12 13 ... + * files (4 time steps per file) = [0 1 2 3][4 5 6 7][8 9 10 11][12 13 ... + * nextFieldData() (1st call) = [0 1] + * nextFieldData() (2nd call) = [1 2] + * nextFieldData() (3rd call) = [2 3] + * nextFieldData() (4th call) = [3 4] + * nextFieldData() (5th call) = [4 5] + * etc... + * + * When getting values from a FieldData using the getVal method, + * the time index is relative to the start of the time-slice. + * + * This means that for d = nextFieldData() (4th call), + * getVal(.., fieldData=d, ..., timeInd = 1, ...) gives a value at + * absolute time step 5 as seen above. + */ FieldData nextFieldData(); ~GPUBufferHandler(); + /** + * You can get the FieldMetaData from here. + */ FieldMetadata *fmd; - static constexpr size_t numberOfTimeStepsPerField = 2; // TODO: Move this to fielddata - static void freeFieldData(); private: diff --git a/src/main.cu b/src/main.cu index 56e3c64..ddc9569 100644 --- a/src/main.cu +++ b/src/main.cu @@ -41,7 +41,6 @@ int main() { middleOfTwoValues<<<1, 1>>>(ptr_test_read, *bufferHandler.fmd, fd); cudaDeviceSynchronize(); - std::cout << "ptr_test_read = " << std::fixed << std::setprecision(6) << *ptr_test_read << "\n"; }