diff --git a/CMakeLists.txt b/CMakeLists.txt index c9f3f275..e0a539fb 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -1,4 +1,4 @@ -project(ScatterAlloc) +project(mallocMC) cmake_minimum_required(VERSION 2.8.5) # helper for libs and packages @@ -48,29 +48,25 @@ INSTALL( DIRECTORY "${CMAKE_CURRENT_SOURCE_DIR}/src/include/." DESTINATION include PATTERN ".git" EXCLUDE - PATTERN "policy_malloc_config.hpp" EXCLUDE + PATTERN "mallocMC_config.hpp" EXCLUDE ) ############################################################################### # Executables ############################################################################### -add_custom_target(examples DEPENDS PolicyAllocExample01 PolicyAllocExample02 VerifyHeap) +add_custom_target(examples DEPENDS mallocMC_Example01 mallocMC_Example02 VerifyHeap) -cuda_add_executable(PolicyAllocExample +cuda_add_executable(mallocMC_Example01 EXCLUDE_FROM_ALL - examples/policy_example.cu ) -cuda_add_executable(PolicyAllocExample01 + examples/mallocMC_example01.cu ) +cuda_add_executable(mallocMC_Example02 EXCLUDE_FROM_ALL - examples/policy_example01.cu ) -cuda_add_executable(PolicyAllocExample02 - EXCLUDE_FROM_ALL - examples/policy_example02.cu ) + examples/mallocMC_example02.cu ) cuda_add_executable(VerifyHeap EXCLUDE_FROM_ALL tests/verify_heap.cu ) -target_link_libraries(PolicyAllocExample ${LIBS}) -target_link_libraries(PolicyAllocExample01 ${LIBS}) -target_link_libraries(PolicyAllocExample02 ${LIBS}) +target_link_libraries(mallocMC_Example01 ${LIBS}) +target_link_libraries(mallocMC_Example02 ${LIBS}) target_link_libraries(VerifyHeap ${LIBS}) diff --git a/INSTALL.md b/INSTALL.md index c76a68fd..669adbd4 100644 --- a/INSTALL.md +++ b/INSTALL.md @@ -29,19 +29,19 @@ Install - `BOOST_ROOT`: Boost installation directory, e.g. `export BOOST_ROOT=` ### Examples -This is an example how to compile `scatteralloc` and test the example code snippets +This is an example how to compile `mallocMC` and test the example code snippets 1. **Setup directories:** - - `mkdir ~/scatteralloc ~/scatterallocBuild` + - `mkdir -p build` 2. **Download the source code:** - - `git clone https://github.com/ComputationalRadiationPhysics/scatteralloc.git ~/scatteralloc` + - `git clone https://github.com/ComputationalRadiationPhysics/mallocMC.git` 3. **Build** - - `cd ~/scatterallocBuild` - - `cmake ../scatteralloc -DCMAKE_INSTALL_PREFIX=$HOME/libs` - - `make` + - `cd build` + - `cmake ../mallocMC -DCMAKE_INSTALL_PREFIX=$HOME/libs` + - `make examples` - `make install` (optional) 4. **Run the examples** - - `./PolicyAllocExample01` - - `./PolicyAllocExample02` + - `./mallocMC_Example01` + - `./mallocMC_Example02` - `./VerifyHeap` - additional options: see `./VerifyHeap --help` diff --git a/LICENSE b/LICENSE index 53fa92a9..667ad8d6 100644 --- a/LICENSE +++ b/LICENSE @@ -1,5 +1,5 @@ /* - ScatterAlloc: Massively Parallel Dynamic Memory Allocation for the GPU. + mallocMC: Memory Allocation for Many Core Architecture http://www.icg.tugraz.at/project/mvp https://www.hzdr.de/crp diff --git a/README.md b/README.md index dc5beb43..82a0079d 100644 --- a/README.md +++ b/README.md @@ -1,10 +1,11 @@ -ScatterAlloc -============ +mallocMC +============= -ScatterAlloc: Massively Parallel Dynamic Memory Allocation for the GPU +mallocMC: Memory Allocator for Many Core Architectures -This project provides a **fast memory manager** for **Nvidia GPUs** with -compute capability `sm_20` or higher. +This project provides a framework for **fast memory managers** on **many core +accelerators**. Currently, it supports **NVIDIA GPUs** of compute capability +`sm_20` or higher through the ScatterAlloc algorithm. From http://www.icg.tugraz.at/project/mvp/downloads : ```quote @@ -28,7 +29,7 @@ Our Homepage: https://www.hzdr.de/crp About This Repository --------------------- -This repository is a +The currently implemented algorithm is a [fork](https://en.wikipedia.org/wiki/Fork_%28software_development%29) of the **ScatterAlloc** project from the [Managed Volume Processing](http://www.icg.tugraz.at/project/mvp) diff --git a/Usage.md b/Usage.md index 2f32d4bf..f9020bbd 100644 --- a/Usage.md +++ b/Usage.md @@ -7,7 +7,7 @@ Step 1: include There is one header file that will include *all* necessary files: ```c++ -#include +#include ``` Step 2a: choose policies @@ -44,7 +44,7 @@ to the policy class: ```c++ // configure the AlignmentPolicy "Shrink" -struct ShrinkConfig : PolicyMalloc::AlignmentPolicies::Shrink<>::Properties { +struct ShrinkConfig : mallocMC::AlignmentPolicies::Shrink<>::Properties { typedef boost::mpl::int_<16> dataAlignment; }; ``` @@ -55,9 +55,9 @@ After configuring the chosen policies, they can be used as template parameters to create the desired allocator type: ```c++ -using namespace PolicyMalloc; +using namespace mallocMC; -typedef PolicyAllocator< +typedef mallocMC::Allocator< CreationPolicy::OldMalloc, DistributionPolicy::Noop, OOMPolicy::ReturnNull, @@ -71,9 +71,9 @@ from NVIDIA CUDA since compute capability sm_20. To get a more novel allocator, could create the following typedef instead: ```c++ -using namespace PolicyMalloc; +using namespace mallocMC; -typedef PolicyAllocator< +typedef mallocMC::Allocator< CreationPolicies::Scatter<>, DistributionPolicies::XMallocSIMD<>, OOMPolicies::ReturnNull, @@ -97,7 +97,7 @@ functions, the following Macro has to be executed: POLICYMALLOC_SET_ALLOCATOR_TYPE(ScatterAllocator) ``` -This will set up the following functions in the namespace `PolicyMalloc`: +This will set up the following functions in the namespace `mallocMC`: | Name | description | |-----------------------|------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------| @@ -105,7 +105,7 @@ This will set up the following functions in the namespace `PolicyMalloc`: | finalizeHeap() | Destroys the heap again | | pbMalloc() / malloc() | Allocates memory on the accelerator | | pbFree() / free() | Frees memory on the accelerator | -| getAvailableSlots() | Determines number of allocatable slots of a certain size. This only works, if the chose CreationPolicy supports it (can be found through `PolicyMalloc::Traits::providesAvailableSlots`) | +| getAvailableSlots() | Determines number of allocatable slots of a certain size. This only works, if the chose CreationPolicy supports it (can be found through `mallocMC::Traits::providesAvailableSlots`) | If the policy class `OldMalloc` is **not** used, it is also possible to execute the Macro @@ -121,37 +121,37 @@ Step 4: use dynamic memory allocation ------------------------------------- A simplistic example would look like this: ```c++ -#include +#include -namespace PolicyMalloc = PM; +namespace mallocMC = MC; -typedef PM::PolicyAllocator< - PM::CreationPolicies::Scatter<>, - PM::DistributionPolicies::XMallocSIMD<>, - PM::OOMPolicies::ReturnNull, - PM::ReservePoolPolicies::SimpleCudaMalloc, - PM::AlignmentPolicies::Shrink +typedef MC::Allocator< + MC::CreationPolicies::Scatter<>, + MC::DistributionPolicies::XMallocSIMD<>, + MC::OOMPolicies::ReturnNull, + MC::ReservePoolPolicies::SimpleCudaMalloc, + MC::AlignmentPolicies::Shrink > ScatterAllocator; -POLICYMALLOC_SET_ALLOCATOR_TYPE(ScatterAllocator) +MALLOCMC_SET_ALLOCATOR_TYPE(ScatterAllocator) __global__ exampleKernel() { // some code ... - int* a = (int*) PM::malloc(sizeof(int)*42); + int* a = (int*) MC::malloc(sizeof(int)*42); // some more code, using *a - PM::free(a); + MC::free(a); } int main(){ - PM::initHeap(512); // heapsize of 512MB + MC::initHeap(512); // heapsize of 512MB exampleKernel<<<32,32>>>(); - PM::finalizeHeap(); + MC::finalizeHeap(); return 0; } ``` diff --git a/examples/example.cu b/examples/example.cu deleted file mode 100644 index 245ea12c..00000000 --- a/examples/example.cu +++ /dev/null @@ -1,132 +0,0 @@ -/* - ScatterAlloc: Massively Parallel Dynamic Memory Allocation for the GPU. - http://www.icg.tugraz.at/project/mvp - - Copyright (C) 2012 Institute for Computer Graphics and Vision, - Graz University of Technology - - Author(s): Markus Steinberger - steinberger ( at ) icg.tugraz.at - Michael Kenzel - kenzel ( at ) icg.tugraz.at - - Permission is hereby granted, free of charge, to any person obtaining a copy - of this software and associated documentation files (the "Software"), to deal - in the Software without restriction, including without limitation the rights - to use, copy, modify, merge, publish, distribute, sublicense, and/or sell - copies of the Software, and to permit persons to whom the Software is - furnished to do so, subject to the following conditions: - - The above copyright notice and this permission notice shall be included in - all copies or substantial portions of the Software. - - THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE - AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER - LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, - OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN - THE SOFTWARE. -*/ - -#include - -//replace the cuda malloc and free calls -#define SCATTERALLOC_OVERWRITE_MALLOC 1 - -//set the template arguments using SCATTERALLOC_HEAPARGS -// pagesize ... byter per page -// accessblocks ... number of superblocks -// regionsize ... number of regions for meta data structur -// wastefactor ... how much memory can be wasted per alloc (multiplicative factor) -// use_coalescing ... combine memory requests of within each warp -// resetfreedpages ... allow pages to be reused with a different size -#define SCATTERALLOC_HEAPARGS 4096, 8, 16, 2, true, false - -//include the scatter alloc heap -#include -#include - -#ifdef WIN32 -#define WIN32_LEAN_AND_MEAN -#define NOMINMAX -#include -#include -#endif - -#include -#include - - -void runexample(int cuda_device); - - -int main(int argc, char** argv) -{ - try - { - int cuda_device = argc > 1 ? atoi(argv[1]) : 0; - - cudaDeviceProp deviceProp; - cudaGetDeviceProperties(&deviceProp, cuda_device); - std::cout << "Using device: " << deviceProp.name << std::endl; - - if( deviceProp.major < 2 ) { - std::cerr << "This GPU with Compute Capability " << deviceProp.major - << "." << deviceProp.minor << " does not meet minimum requirements." << std::endl; - std::cerr << "A GPU with Compute Capability >= 2.0 is required." << std::endl; - return -2; - } - - runexample(cuda_device); - - cudaDeviceReset(); - } - catch (const std::exception& e) - { - std::cout << e.what() << std::endl; - #ifdef WIN32 - while (!_kbhit()); - #endif - return -1; - } - catch (...) - { - std::cout << "unknown exception!" << std::endl; - #ifdef WIN32 - while (!_kbhit()); - #endif - return -1; - } - - return 0; -} - - -__global__ void allocSomething(uint** parray) -{ - parray[threadIdx.x + blockIdx.x*blockDim.x] = new uint[threadIdx.x % 4]; -} -__global__ void freeSomething(uint** parray) -{ - delete[] parray[threadIdx.x + blockIdx.x*blockDim.x]; -} - - -void runexample(int cuda_device) -{ - cudaSetDevice(cuda_device); - - //init the heap - initHeap(); - //you can also specify the size of the heap in bytes - //initHeap(8U*1024U*1024U); - - size_t block = 128; - size_t grid = 64; - - uint** data; - SCATTERALLOC_CUDA_CHECKED_CALL(cudaMalloc(&data, grid*block*sizeof(uint*))); - allocSomething<<>>(data); - SCATTERALLOC_CUDA_CHECKED_CALL(cudaDeviceSynchronize()); - freeSomething<<>>(data); - SCATTERALLOC_CUDA_CHECKED_CALL(cudaDeviceSynchronize()); -} diff --git a/examples/policy_example01.cu b/examples/mallocMC_example01.cu similarity index 82% rename from examples/policy_example01.cu rename to examples/mallocMC_example01.cu index 8aae8da6..7cf5d3dc 100644 --- a/examples/policy_example01.cu +++ b/examples/mallocMC_example01.cu @@ -1,5 +1,5 @@ /* - ScatterAlloc: Massively Parallel Dynamic Memory Allocation for the GPU. + mallocMC: Memory Allocator for Many Core Architectures. https://www.hzdr.de/crp Copyright 2014 Institute of Radiation Physics, @@ -32,7 +32,7 @@ #include #include -#include "policy_example01_config.hpp" +#include "mallocMC_example01_config.cu" void run(); @@ -61,18 +61,18 @@ __device__ int** c; __global__ void createArrays(int x, int y){ - a = (int**) PolicyMalloc::malloc(sizeof(int*) * x*y); - b = (int**) PolicyMalloc::malloc(sizeof(int*) * x*y); - c = (int**) PolicyMalloc::malloc(sizeof(int*) * x*y); + a = (int**) mallocMC::malloc(sizeof(int*) * x*y); + b = (int**) mallocMC::malloc(sizeof(int*) * x*y); + c = (int**) mallocMC::malloc(sizeof(int*) * x*y); } __global__ void fillArrays(int length, int* d){ int id = threadIdx.x + blockIdx.x*blockDim.x; - a[id] = (int*) PolicyMalloc::malloc(length*sizeof(int)); - b[id] = (int*) PolicyMalloc::malloc(length*sizeof(int)); - c[id] = (int*) PolicyMalloc::malloc(sizeof(int)*length); + a[id] = (int*) mallocMC::malloc(length*sizeof(int)); + b[id] = (int*) mallocMC::malloc(length*sizeof(int)); + c[id] = (int*) mallocMC::malloc(sizeof(int)*length); for(int i=0 ; i>>(); cudaFree(d); //finalize the heap again - PolicyMalloc::finalizeHeap(); + mallocMC::finalizeHeap(); } diff --git a/examples/policy_example01_config.hpp b/examples/mallocMC_example01_config.cu similarity index 63% rename from examples/policy_example01_config.hpp rename to examples/mallocMC_example01_config.cu index cb81530c..1c271137 100644 --- a/examples/policy_example01_config.hpp +++ b/examples/mallocMC_example01_config.cu @@ -1,5 +1,5 @@ /* - ScatterAlloc: Massively Parallel Dynamic Memory Allocation for the GPU. + mallocMC: Memory Allocator for Many Core Architectures. https://www.hzdr.de/crp Copyright 2014 Institute of Radiation Physics, @@ -31,21 +31,21 @@ #include #include -// basic files for PolicyMalloc -#include "src/include/scatteralloc/policy_malloc_overwrites.hpp" -#include "src/include/scatteralloc/policy_malloc_hostclass.hpp" +// basic files for mallocMC +#include "src/include/mallocMC/mallocMC_overwrites.hpp" +#include "src/include/mallocMC/mallocMC_hostclass.hpp" -// Load all available policies for PolicyMalloc -#include "src/include/scatteralloc/CreationPolicies.hpp" -#include "src/include/scatteralloc/DistributionPolicies.hpp" -#include "src/include/scatteralloc/OOMPolicies.hpp" -#include "src/include/scatteralloc/ReservePoolPolicies.hpp" -#include "src/include/scatteralloc/AlignmentPolicies.hpp" +// Load all available policies for mallocMC +#include "src/include/mallocMC/CreationPolicies.hpp" +#include "src/include/mallocMC/DistributionPolicies.hpp" +#include "src/include/mallocMC/OOMPolicies.hpp" +#include "src/include/mallocMC/ReservePoolPolicies.hpp" +#include "src/include/mallocMC/AlignmentPolicies.hpp" // configurate the CreationPolicy "Scatter" to modify the default behaviour -struct ScatterHeapConfig : PolicyMalloc::CreationPolicies::Scatter<>::HeapProperties{ +struct ScatterHeapConfig : mallocMC::CreationPolicies::Scatter<>::HeapProperties{ typedef boost::mpl::int_<4096> pagesize; typedef boost::mpl::int_<8> accessblocks; typedef boost::mpl::int_<16> regionsize; @@ -53,7 +53,7 @@ struct ScatterHeapConfig : PolicyMalloc::CreationPolicies::Scatter<>::HeapProper typedef boost::mpl::bool_ resetfreedpages; }; -struct ScatterHashConfig : PolicyMalloc::CreationPolicies::Scatter<>::HashingProperties{ +struct ScatterHashConfig : mallocMC::CreationPolicies::Scatter<>::HashingProperties{ typedef boost::mpl::int_<38183> hashingK; typedef boost::mpl::int_<17497> hashingDistMP; typedef boost::mpl::int_<1> hashingDistWP; @@ -61,24 +61,24 @@ struct ScatterHashConfig : PolicyMalloc::CreationPolicies::Scatter<>::HashingPro }; // configure the DistributionPolicy "XMallocSIMD" -struct XMallocConfig : PolicyMalloc::DistributionPolicies::XMallocSIMD<>::Properties { +struct XMallocConfig : mallocMC::DistributionPolicies::XMallocSIMD<>::Properties { typedef ScatterHeapConfig::pagesize pagesize; }; // configure the AlignmentPolicy "Shrink" -struct ShrinkConfig : PolicyMalloc::AlignmentPolicies::Shrink<>::Properties { +struct ShrinkConfig : mallocMC::AlignmentPolicies::Shrink<>::Properties { typedef boost::mpl::int_<16> dataAlignment; }; // Define a new allocator and call it ScatterAllocator // which resembles the behaviour of ScatterAlloc -typedef PolicyMalloc::PolicyAllocator< - PolicyMalloc::CreationPolicies::Scatter, - PolicyMalloc::DistributionPolicies::XMallocSIMD, - PolicyMalloc::OOMPolicies::ReturnNull, - PolicyMalloc::ReservePoolPolicies::SimpleCudaMalloc, - PolicyMalloc::AlignmentPolicies::Shrink +typedef mallocMC::Allocator< + mallocMC::CreationPolicies::Scatter, + mallocMC::DistributionPolicies::XMallocSIMD, + mallocMC::OOMPolicies::ReturnNull, + mallocMC::ReservePoolPolicies::SimpleCudaMalloc, + mallocMC::AlignmentPolicies::Shrink > ScatterAllocator; // use "ScatterAllocator" as PolicyAllocator -POLICYMALLOC_SET_ALLOCATOR_TYPE(ScatterAllocator) +MALLOCMC_SET_ALLOCATOR_TYPE(ScatterAllocator) diff --git a/examples/policy_example02.cu b/examples/mallocMC_example02.cu similarity index 77% rename from examples/policy_example02.cu rename to examples/mallocMC_example02.cu index ea7266a4..167c1bed 100644 --- a/examples/policy_example02.cu +++ b/examples/mallocMC_example02.cu @@ -1,5 +1,5 @@ /* - ScatterAlloc: Massively Parallel Dynamic Memory Allocation for the GPU. + mallocMC: Memory Allocator for Many Core Architectures. https://www.hzdr.de/crp Copyright 2014 Institute of Radiation Physics, @@ -36,21 +36,21 @@ #include /////////////////////////////////////////////////////////////////////////////// -// includes for PolicyMalloc +// includes for mallocMC /////////////////////////////////////////////////////////////////////////////// -// basic files for PolicyMalloc -#include "src/include/scatteralloc/policy_malloc_overwrites.hpp" -#include "src/include/scatteralloc/policy_malloc_hostclass.hpp" - -// Load all available policies for PolicyMalloc -#include "src/include/scatteralloc/CreationPolicies.hpp" -#include "src/include/scatteralloc/DistributionPolicies.hpp" -#include "src/include/scatteralloc/OOMPolicies.hpp" -#include "src/include/scatteralloc/ReservePoolPolicies.hpp" -#include "src/include/scatteralloc/AlignmentPolicies.hpp" +// basic files for mallocMC +#include "src/include/mallocMC/mallocMC_overwrites.hpp" +#include "src/include/mallocMC/mallocMC_hostclass.hpp" + +// Load all available policies for mallocMC +#include "src/include/mallocMC/CreationPolicies.hpp" +#include "src/include/mallocMC/DistributionPolicies.hpp" +#include "src/include/mallocMC/OOMPolicies.hpp" +#include "src/include/mallocMC/ReservePoolPolicies.hpp" +#include "src/include/mallocMC/AlignmentPolicies.hpp" /////////////////////////////////////////////////////////////////////////////// -// Configuration for PolicyMalloc +// Configuration for mallocMC /////////////////////////////////////////////////////////////////////////////// // configurate the CreationPolicy "Scatter" @@ -81,23 +81,23 @@ struct AlignmentConfig{ // Define a new allocator and call it ScatterAllocator // which resembles the behaviour of ScatterAlloc -typedef PolicyMalloc::PolicyAllocator< - PolicyMalloc::CreationPolicies::Scatter, - PolicyMalloc::DistributionPolicies::XMallocSIMD, - PolicyMalloc::OOMPolicies::ReturnNull, - PolicyMalloc::ReservePoolPolicies::SimpleCudaMalloc, - PolicyMalloc::AlignmentPolicies::Shrink +typedef mallocMC::Allocator< + mallocMC::CreationPolicies::Scatter, + mallocMC::DistributionPolicies::XMallocSIMD, + mallocMC::OOMPolicies::ReturnNull, + mallocMC::ReservePoolPolicies::SimpleCudaMalloc, + mallocMC::AlignmentPolicies::Shrink > ScatterAllocator; -// use "ScatterAllocator" as PolicyAllocator -POLICYMALLOC_SET_ALLOCATOR_TYPE(ScatterAllocator) +// use "ScatterAllocator" as mallocMC +MALLOCMC_SET_ALLOCATOR_TYPE(ScatterAllocator) -// replace all standard malloc()-calls on the device by PolicyAllocator calls +// replace all standard malloc()-calls on the device by mallocMC calls // This will not work with the CreationPolicy "OldMalloc"! -POLICYMALLOC_OVERWRITE_MALLOC() +MALLOCMC_OVERWRITE_MALLOC() /////////////////////////////////////////////////////////////////////////////// -// End of PolicyMalloc configuration +// End of mallocMC configuration /////////////////////////////////////////////////////////////////////////////// @@ -137,7 +137,7 @@ __global__ void createArrays(int x, int y){ __global__ void fillArrays(int length, int* d){ int id = threadIdx.x + blockIdx.x*blockDim.x; - // using the POLICYMALLOC_OVERWRITE_MALLOC() macro + // using the MALLOCMC_OVERWRITE_MALLOC() macro // allows also the use of "new" a[id] = new int[length]; b[id] = new int[length]; @@ -178,7 +178,7 @@ void run() //init the heap std::cerr << "initHeap..."; - PolicyMalloc::initHeap(1U*1024U*1024U*1024U); //1GB for device-side malloc + mallocMC::initHeap(1U*1024U*1024U*1024U); //1GB for device-side malloc std::cerr << "done" << std::endl; // device-side pointers @@ -207,14 +207,14 @@ void run() int gaussian = n*(n-1); std::cout << "The gaussian sum as comparison: " << gaussian << std::endl; - if(PolicyMalloc::Traits::providesAvailableSlots){ + if(mallocMC::Traits::providesAvailableSlots){ std::cout << "there are "; - std::cout << PolicyMalloc::getAvailableSlots(1024U*1024U); + std::cout << mallocMC::getAvailableSlots(1024U*1024U); std::cout << " Slots of size 1MB available" << std::endl; } freeArrays<<>>(); cudaFree(d); //finalize the heap again - PolicyMalloc::finalizeHeap(); + mallocMC::finalizeHeap(); } diff --git a/examples/policy_example.cu b/examples/policy_example.cu deleted file mode 100644 index 99580ad5..00000000 --- a/examples/policy_example.cu +++ /dev/null @@ -1,133 +0,0 @@ -/* - ScatterAlloc: Massively Parallel Dynamic Memory Allocation for the GPU. - http://www.icg.tugraz.at/project/mvp - - Copyright (C) 2012 Institute for Computer Graphics and Vision, - Graz University of Technology - - Author(s): Markus Steinberger - steinberger ( at ) icg.tugraz.at - Michael Kenzel - kenzel ( at ) icg.tugraz.at - - Permission is hereby granted, free of charge, to any person obtaining a copy - of this software and associated documentation files (the "Software"), to deal - in the Software without restriction, including without limitation the rights - to use, copy, modify, merge, publish, distribute, sublicense, and/or sell - copies of the Software, and to permit persons to whom the Software is - furnished to do so, subject to the following conditions: - - The above copyright notice and this permission notice shall be included in - all copies or substantial portions of the Software. - - THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE - AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER - LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, - OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN - THE SOFTWARE. -*/ - -#include -#include "src/include/scatteralloc/policy_malloc_utils.hpp" - -#include "src/include/scatteralloc/policy_malloc_config.hpp" - -using PolicyMalloc::initHeap; -#ifdef __CUDACC__ -#if __CUDA_ARCH >= 200 -using PolicyMalloc::pbMalloc; -using PolicyMalloc::pbFree; -#endif -#endif - - - - -#ifdef WIN32 -#define WIN32_LEAN_AND_MEAN -#define NOMINMAX -#include -#include -#endif - -#include -#include - -typedef PolicyMalloc::uint32_richtig_huebsch uint; - -void runexample(int cuda_device); - - -int main(int argc, char** argv) -{ - try - { - int cuda_device = argc > 1 ? atoi(argv[1]) : 0; - - cudaDeviceProp deviceProp; - cudaGetDeviceProperties(&deviceProp, cuda_device); - std::cout << "Using device: " << deviceProp.name << std::endl; - - if( deviceProp.major < 2 ) { - std::cerr << "This GPU with Compute Capability " << deviceProp.major - << "." << deviceProp.minor << " does not meet minimum requirements." << std::endl; - std::cerr << "A GPU with Compute Capability >= 2.0 is required." << std::endl; - return -2; - } - - std::cout << "start" <>>(data); - SCATTERALLOC_CUDA_CHECKED_CALL(cudaDeviceSynchronize()); - freeSomething<<>>(data); - SCATTERALLOC_CUDA_CHECKED_CALL(cudaDeviceSynchronize()); -} diff --git a/src/include/scatteralloc/AlignmentPolicies.hpp b/src/include/mallocMC/AlignmentPolicies.hpp similarity index 95% rename from src/include/scatteralloc/AlignmentPolicies.hpp rename to src/include/mallocMC/AlignmentPolicies.hpp index af874d18..c471b696 100644 --- a/src/include/scatteralloc/AlignmentPolicies.hpp +++ b/src/include/mallocMC/AlignmentPolicies.hpp @@ -1,5 +1,5 @@ /* - ScatterAlloc: Massively Parallel Dynamic Memory Allocation for the GPU. + mallocMC: Memory Allocator for Many Core Architectures. https://www.hzdr.de/crp Copyright 2014 Institute of Radiation Physics, diff --git a/src/include/scatteralloc/CreationPolicies.hpp b/src/include/mallocMC/CreationPolicies.hpp similarity index 95% rename from src/include/scatteralloc/CreationPolicies.hpp rename to src/include/mallocMC/CreationPolicies.hpp index d037d456..56f6e23f 100644 --- a/src/include/scatteralloc/CreationPolicies.hpp +++ b/src/include/mallocMC/CreationPolicies.hpp @@ -1,5 +1,5 @@ /* - ScatterAlloc: Massively Parallel Dynamic Memory Allocation for the GPU. + mallocMC: Memory Allocator for Many Core Architectures. https://www.hzdr.de/crp Copyright 2014 Institute of Radiation Physics, diff --git a/src/include/scatteralloc/DistributionPolicies.hpp b/src/include/mallocMC/DistributionPolicies.hpp similarity index 95% rename from src/include/scatteralloc/DistributionPolicies.hpp rename to src/include/mallocMC/DistributionPolicies.hpp index 2bd32ab7..f534c57d 100644 --- a/src/include/scatteralloc/DistributionPolicies.hpp +++ b/src/include/mallocMC/DistributionPolicies.hpp @@ -1,5 +1,5 @@ /* - ScatterAlloc: Massively Parallel Dynamic Memory Allocation for the GPU. + mallocMC: Memory Allocator for Many Core Architectures. https://www.hzdr.de/crp Copyright 2014 Institute of Radiation Physics, diff --git a/src/include/scatteralloc/OOMPolicies.hpp b/src/include/mallocMC/OOMPolicies.hpp similarity index 95% rename from src/include/scatteralloc/OOMPolicies.hpp rename to src/include/mallocMC/OOMPolicies.hpp index 11c8a546..67eda519 100644 --- a/src/include/scatteralloc/OOMPolicies.hpp +++ b/src/include/mallocMC/OOMPolicies.hpp @@ -1,5 +1,5 @@ /* - ScatterAlloc: Massively Parallel Dynamic Memory Allocation for the GPU. + mallocMC: Memory Allocator for Many Core Architectures. https://www.hzdr.de/crp Copyright 2014 Institute of Radiation Physics, diff --git a/src/include/scatteralloc/ReservePoolPolicies.hpp b/src/include/mallocMC/ReservePoolPolicies.hpp similarity index 95% rename from src/include/scatteralloc/ReservePoolPolicies.hpp rename to src/include/mallocMC/ReservePoolPolicies.hpp index 34848506..1cc3aa7b 100644 --- a/src/include/scatteralloc/ReservePoolPolicies.hpp +++ b/src/include/mallocMC/ReservePoolPolicies.hpp @@ -1,5 +1,5 @@ /* - ScatterAlloc: Massively Parallel Dynamic Memory Allocation for the GPU. + mallocMC: Memory Allocator for Many Core Architectures. https://www.hzdr.de/crp Copyright 2014 Institute of Radiation Physics, diff --git a/src/include/scatteralloc/alignmentPolicies/Noop.hpp b/src/include/mallocMC/alignmentPolicies/Noop.hpp similarity index 92% rename from src/include/scatteralloc/alignmentPolicies/Noop.hpp rename to src/include/mallocMC/alignmentPolicies/Noop.hpp index 456d71ce..3879da6f 100644 --- a/src/include/scatteralloc/alignmentPolicies/Noop.hpp +++ b/src/include/mallocMC/alignmentPolicies/Noop.hpp @@ -1,5 +1,5 @@ /* - ScatterAlloc: Massively Parallel Dynamic Memory Allocation for the GPU. + mallocMC: Memory Allocator for Many Core Architectures. Copyright 2014 Institute of Radiation Physics, Helmholtz-Zentrum Dresden - Rossendorf @@ -27,7 +27,7 @@ #pragma once -namespace PolicyMalloc{ +namespace mallocMC{ namespace AlignmentPolicies{ /** @@ -39,4 +39,4 @@ namespace AlignmentPolicies{ class Noop; } //namespace AlignmentPolicies -} //namespace PolicyMalloc +} //namespace mallocMC diff --git a/src/include/scatteralloc/alignmentPolicies/Noop_impl.hpp b/src/include/mallocMC/alignmentPolicies/Noop_impl.hpp similarity index 90% rename from src/include/scatteralloc/alignmentPolicies/Noop_impl.hpp rename to src/include/mallocMC/alignmentPolicies/Noop_impl.hpp index 1a454748..e1c034d0 100644 --- a/src/include/scatteralloc/alignmentPolicies/Noop_impl.hpp +++ b/src/include/mallocMC/alignmentPolicies/Noop_impl.hpp @@ -1,5 +1,5 @@ /* - ScatterAlloc: Massively Parallel Dynamic Memory Allocation for the GPU. + mallocMC: Memory Allocator for Many Core Architectures. Copyright 2014 Institute of Radiation Physics, Helmholtz-Zentrum Dresden - Rossendorf @@ -31,9 +31,9 @@ #include #include "Noop.hpp" -#include "../policy_malloc_prefixes.hpp" +#include "../mallocMC_prefixes.hpp" -namespace PolicyMalloc{ +namespace mallocMC{ namespace AlignmentPolicies{ class Noop{ @@ -45,7 +45,7 @@ namespace AlignmentPolicies{ return boost::make_tuple(memory,memsize); } - PMMA_ACCELERATOR + MAMC_HOST MAMC_ACCELERATOR static uint32 applyPadding(uint32 bytes){ return bytes; } @@ -57,4 +57,4 @@ namespace AlignmentPolicies{ }; } //namespace AlignmentPolicies -} //namespace PolicyMalloc +} //namespace mallocMC diff --git a/src/include/scatteralloc/alignmentPolicies/Shrink.hpp b/src/include/mallocMC/alignmentPolicies/Shrink.hpp similarity index 94% rename from src/include/scatteralloc/alignmentPolicies/Shrink.hpp rename to src/include/mallocMC/alignmentPolicies/Shrink.hpp index eab7169a..4bafeff8 100644 --- a/src/include/scatteralloc/alignmentPolicies/Shrink.hpp +++ b/src/include/mallocMC/alignmentPolicies/Shrink.hpp @@ -1,5 +1,5 @@ /* - ScatterAlloc: Massively Parallel Dynamic Memory Allocation for the GPU. + mallocMC: Memory Allocator for Many Core Architectures. http://www.icg.tugraz.at/project/mvp Copyright (C) 2012 Institute for Computer Graphics and Vision, @@ -33,7 +33,7 @@ #include -namespace PolicyMalloc{ +namespace mallocMC{ namespace AlignmentPolicies{ namespace ShrinkConfig{ @@ -58,4 +58,4 @@ namespace ShrinkConfig{ class Shrink; } //namespace AlignmentPolicies -} //namespace PolicyMalloc +} //namespace mallocMC diff --git a/src/include/scatteralloc/alignmentPolicies/Shrink_impl.hpp b/src/include/mallocMC/alignmentPolicies/Shrink_impl.hpp similarity index 89% rename from src/include/scatteralloc/alignmentPolicies/Shrink_impl.hpp rename to src/include/mallocMC/alignmentPolicies/Shrink_impl.hpp index fc476346..dd80ddd4 100644 --- a/src/include/scatteralloc/alignmentPolicies/Shrink_impl.hpp +++ b/src/include/mallocMC/alignmentPolicies/Shrink_impl.hpp @@ -1,5 +1,5 @@ /* - ScatterAlloc: Massively Parallel Dynamic Memory Allocation for the GPU. + mallocMC: Memory Allocator for Many Core Architectures. http://www.icg.tugraz.at/project/mvp Copyright (C) 2012 Institute for Computer Graphics and Vision, @@ -39,9 +39,9 @@ #include #include "Shrink.hpp" -#include "../policy_malloc_prefixes.hpp" +#include "../mallocMC_prefixes.hpp" -namespace PolicyMalloc{ +namespace mallocMC{ namespace AlignmentPolicies{ namespace Shrink2NS{ @@ -65,14 +65,14 @@ namespace Shrink2NS{ * shipped default-parameters (in the inherited struct) have lowest precedence. * They will be overridden by a given configuration struct. However, even the * given configuration struct can be overridden by compile-time command line - * parameters (e.g. -D POLICYMALLOC_AP_SHRINK_DATAALIGNMENT 128) + * parameters (e.g. -D MALLOCMC_AP_SHRINK_DATAALIGNMENT 128) * * default-struct < template-struct < command-line parameter */ -#ifndef POLICYMALLOC_AP_SHRINK_DATAALIGNMENT -#define POLICYMALLOC_AP_SHRINK_DATAALIGNMENT Properties::dataAlignment::value +#ifndef MALLOCMC_AP_SHRINK_DATAALIGNMENT +#define MALLOCMC_AP_SHRINK_DATAALIGNMENT Properties::dataAlignment::value #endif - static const uint32 dataAlignment = POLICYMALLOC_AP_SHRINK_DATAALIGNMENT; + static const uint32 dataAlignment = MALLOCMC_AP_SHRINK_DATAALIGNMENT; BOOST_STATIC_ASSERT(dataAlignment > 0); //dataAlignment must also be a power of 2! @@ -102,13 +102,13 @@ namespace Shrink2NS{ return boost::make_tuple(memory,memsize); } - PMMA_HOST - PMMA_ACCELERATOR + MAMC_HOST + MAMC_ACCELERATOR static uint32 applyPadding(uint32 bytes){ return (bytes + dataAlignment - 1) & ~(dataAlignment-1); } - PMMA_HOST + MAMC_HOST static std::string classname(){ std::stringstream ss; ss << "Shrink[" << dataAlignment << "]"; @@ -118,4 +118,4 @@ namespace Shrink2NS{ }; } //namespace AlignmentPolicies -} //namespace PolicyMalloc +} //namespace mallocMC diff --git a/src/include/scatteralloc/creationPolicies/OldMalloc.hpp b/src/include/mallocMC/creationPolicies/OldMalloc.hpp similarity index 92% rename from src/include/scatteralloc/creationPolicies/OldMalloc.hpp rename to src/include/mallocMC/creationPolicies/OldMalloc.hpp index b131ed62..e48b4890 100644 --- a/src/include/scatteralloc/creationPolicies/OldMalloc.hpp +++ b/src/include/mallocMC/creationPolicies/OldMalloc.hpp @@ -1,5 +1,5 @@ /* - ScatterAlloc: Massively Parallel Dynamic Memory Allocation for the GPU. + mallocMC: Memory Allocator for Many Core Architectures. Copyright 2014 Institute of Radiation Physics, Helmholtz-Zentrum Dresden - Rossendorf @@ -28,7 +28,7 @@ #pragma once -namespace PolicyMalloc{ +namespace mallocMC{ namespace CreationPolicies{ /** @@ -41,4 +41,4 @@ namespace CreationPolicies{ class OldMalloc; } //namespace CreationPolicies -} //namespace PolicyMalloc +} //namespace mallocMC diff --git a/src/include/scatteralloc/creationPolicies/OldMalloc_impl.hpp b/src/include/mallocMC/creationPolicies/OldMalloc_impl.hpp similarity index 91% rename from src/include/scatteralloc/creationPolicies/OldMalloc_impl.hpp rename to src/include/mallocMC/creationPolicies/OldMalloc_impl.hpp index 970034d3..ab4b66c5 100644 --- a/src/include/scatteralloc/creationPolicies/OldMalloc_impl.hpp +++ b/src/include/mallocMC/creationPolicies/OldMalloc_impl.hpp @@ -1,5 +1,5 @@ /* - ScatterAlloc: Massively Parallel Dynamic Memory Allocation for the GPU. + mallocMC: Memory Allocator for Many Core Architectures. Copyright 2014 Institute of Radiation Physics, Helmholtz-Zentrum Dresden - Rossendorf @@ -32,7 +32,7 @@ #include "OldMalloc.hpp" -namespace PolicyMalloc{ +namespace mallocMC{ namespace CreationPolicies{ class OldMalloc @@ -59,7 +59,7 @@ namespace CreationPolicies{ template < typename T> static void* initHeap(const T& obj, void* pool, size_t memsize){ T* dAlloc; - SCATTERALLOC_CUDA_CHECKED_CALL(cudaGetSymbolAddress((void**)&dAlloc,obj)); + MALLOCMC_CUDA_CHECKED_CALL(cudaGetSymbolAddress((void**)&dAlloc,obj)); return dAlloc; } @@ -75,4 +75,4 @@ namespace CreationPolicies{ }; } //namespace CreationPolicies -} //namespace PolicyMalloc +} //namespace mallocMC diff --git a/src/include/scatteralloc/creationPolicies/Scatter.hpp b/src/include/mallocMC/creationPolicies/Scatter.hpp similarity index 96% rename from src/include/scatteralloc/creationPolicies/Scatter.hpp rename to src/include/mallocMC/creationPolicies/Scatter.hpp index 3f9182b1..499bd73e 100644 --- a/src/include/scatteralloc/creationPolicies/Scatter.hpp +++ b/src/include/mallocMC/creationPolicies/Scatter.hpp @@ -1,5 +1,5 @@ /* - ScatterAlloc: Massively Parallel Dynamic Memory Allocation for the GPU. + mallocMC: Memory Allocator for Many Core Architectures. http://www.icg.tugraz.at/project/mvp Copyright (C) 2012 Institute for Computer Graphics and Vision, @@ -36,7 +36,7 @@ #include #include -namespace PolicyMalloc{ +namespace mallocMC{ namespace CreationPolicies{ namespace ScatterConf{ struct DefaultScatterConfig{ @@ -81,4 +81,4 @@ namespace ScatterConf{ class Scatter; }// namespace CreationPolicies -}// namespace PolicyMalloc +}// namespace mallocMC diff --git a/src/include/scatteralloc/creationPolicies/Scatter_impl.hpp b/src/include/mallocMC/creationPolicies/Scatter_impl.hpp similarity index 93% rename from src/include/scatteralloc/creationPolicies/Scatter_impl.hpp rename to src/include/mallocMC/creationPolicies/Scatter_impl.hpp index f20d2d38..d38871bb 100644 --- a/src/include/scatteralloc/creationPolicies/Scatter_impl.hpp +++ b/src/include/mallocMC/creationPolicies/Scatter_impl.hpp @@ -1,5 +1,5 @@ /* - ScatterAlloc: Massively Parallel Dynamic Memory Allocation for the GPU. + mallocMC: Memory Allocator for Many Core Architectures. http://www.icg.tugraz.at/project/mvp Copyright (C) 2012 Institute for Computer Graphics and Vision, @@ -40,10 +40,10 @@ #include #include -#include "../policy_malloc_utils.hpp" +#include "../mallocMC_utils.hpp" #include "Scatter.hpp" -namespace PolicyMalloc{ +namespace mallocMC{ namespace CreationPolicies{ namespace ScatterKernelDetail{ @@ -88,34 +88,34 @@ namespace ScatterKernelDetail{ * shipped default-parameters (in the inherited struct) have lowest precedence. * They will be overridden by a given configuration struct. However, even the * given configuration struct can be overridden by compile-time command line - * parameters (e.g. -D POLICYMALLOC_CP_SCATTER_PAGESIZE 1024) + * parameters (e.g. -D MALLOCMC_CP_SCATTER_PAGESIZE 1024) * * default-struct < template-struct < command-line parameter */ -#ifndef POLICYMALLOC_CP_SCATTER_PAGESIZE -#define POLICYMALLOC_CP_SCATTER_PAGESIZE static_cast(HeapProperties::pagesize::value) +#ifndef MALLOCMC_CP_SCATTER_PAGESIZE +#define MALLOCMC_CP_SCATTER_PAGESIZE static_cast(HeapProperties::pagesize::value) #endif - static const uint32 pagesize = POLICYMALLOC_CP_SCATTER_PAGESIZE; + static const uint32 pagesize = MALLOCMC_CP_SCATTER_PAGESIZE; -#ifndef POLICYMALLOC_CP_SCATTER_ACCESSBLOCKS -#define POLICYMALLOC_CP_SCATTER_ACCESSBLOCKS static_cast(HeapProperties::accessblocks::value) +#ifndef MALLOCMC_CP_SCATTER_ACCESSBLOCKS +#define MALLOCMC_CP_SCATTER_ACCESSBLOCKS static_cast(HeapProperties::accessblocks::value) #endif - static const uint32 accessblocks = POLICYMALLOC_CP_SCATTER_ACCESSBLOCKS; + static const uint32 accessblocks = MALLOCMC_CP_SCATTER_ACCESSBLOCKS; -#ifndef POLICYMALLOC_CP_SCATTER_REGIONSIZE -#define POLICYMALLOC_CP_SCATTER_REGIONSIZE static_cast(HeapProperties::regionsize::value) +#ifndef MALLOCMC_CP_SCATTER_REGIONSIZE +#define MALLOCMC_CP_SCATTER_REGIONSIZE static_cast(HeapProperties::regionsize::value) #endif - static const uint32 regionsize = POLICYMALLOC_CP_SCATTER_REGIONSIZE; + static const uint32 regionsize = MALLOCMC_CP_SCATTER_REGIONSIZE; -#ifndef POLICYMALLOC_CP_SCATTER_WASTEFACTOR -#define POLICYMALLOC_CP_SCATTER_WASTEFACTOR static_cast(HeapProperties::wastefactor::value) +#ifndef MALLOCMC_CP_SCATTER_WASTEFACTOR +#define MALLOCMC_CP_SCATTER_WASTEFACTOR static_cast(HeapProperties::wastefactor::value) #endif - static const uint32 wastefactor = POLICYMALLOC_CP_SCATTER_WASTEFACTOR; + static const uint32 wastefactor = MALLOCMC_CP_SCATTER_WASTEFACTOR; -#ifndef POLICYMALLOC_CP_SCATTER_RESETFREEDPAGES -#define POLICYMALLOC_CP_SCATTER_RESETFREEDPAGES static_cast(HeapProperties::resetfreedpages::value) +#ifndef MALLOCMC_CP_SCATTER_RESETFREEDPAGES +#define MALLOCMC_CP_SCATTER_RESETFREEDPAGES static_cast(HeapProperties::resetfreedpages::value) #endif - static const bool resetfreedpages = POLICYMALLOC_CP_SCATTER_RESETFREEDPAGES; + static const bool resetfreedpages = MALLOCMC_CP_SCATTER_RESETFREEDPAGES; public: @@ -133,25 +133,25 @@ namespace ScatterKernelDetail{ static const uint32 minChunkSize1 = 0x10; static const uint32 HierarchyThreshold = (pagesize - 2*sizeof(uint32))/33; -#ifndef POLICYMALLOC_CP_SCATTER_HASHINGK -#define POLICYMALLOC_CP_SCATTER_HASHINGK static_cast(HashingProperties::hashingK::value) +#ifndef MALLOCMC_CP_SCATTER_HASHINGK +#define MALLOCMC_CP_SCATTER_HASHINGK static_cast(HashingProperties::hashingK::value) #endif - static const uint32 hashingK = POLICYMALLOC_CP_SCATTER_HASHINGK; + static const uint32 hashingK = MALLOCMC_CP_SCATTER_HASHINGK; -#ifndef POLICYMALLOC_CP_SCATTER_HASHINGDISTMP -#define POLICYMALLOC_CP_SCATTER_HASHINGDISTMP static_cast(HashingProperties::hashingDistMP::value) +#ifndef MALLOCMC_CP_SCATTER_HASHINGDISTMP +#define MALLOCMC_CP_SCATTER_HASHINGDISTMP static_cast(HashingProperties::hashingDistMP::value) #endif - static const uint32 hashingDistMP = POLICYMALLOC_CP_SCATTER_HASHINGDISTMP; + static const uint32 hashingDistMP = MALLOCMC_CP_SCATTER_HASHINGDISTMP; -#ifndef POLICYMALLOC_CP_SCATTER_HASHINGDISTWP -#define POLICYMALLOC_CP_SCATTER_HASHINGDISTWP static_cast(HashingProperties::hashingDistWP::value) +#ifndef MALLOCMC_CP_SCATTER_HASHINGDISTWP +#define MALLOCMC_CP_SCATTER_HASHINGDISTWP static_cast(HashingProperties::hashingDistWP::value) #endif - static const uint32 hashingDistWP = POLICYMALLOC_CP_SCATTER_HASHINGDISTWP; + static const uint32 hashingDistWP = MALLOCMC_CP_SCATTER_HASHINGDISTWP; -#ifndef POLICYMALLOC_CP_SCATTER_HASHINGDISTWPREL -#define POLICYMALLOC_CP_SCATTER_HASHINGDISTWPREL static_cast(HashingProperties::hashingDistWPRel::value) +#ifndef MALLOCMC_CP_SCATTER_HASHINGDISTWPREL +#define MALLOCMC_CP_SCATTER_HASHINGDISTWPREL static_cast(HashingProperties::hashingDistWPRel::value) #endif - static const uint32 hashingDistWPRel = POLICYMALLOC_CP_SCATTER_HASHINGDISTWPREL; + static const uint32 hashingDistWPRel = MALLOCMC_CP_SCATTER_HASHINGDISTWPREL; /** @@ -786,7 +786,7 @@ namespace ScatterKernelDetail{ template < typename T_Obj> static void* initHeap(const T_Obj& obj, void* pool, size_t memsize){ T_Obj* heap; - SCATTERALLOC_CUDA_CHECKED_CALL(cudaGetSymbolAddress((void**)&heap,obj)); + MALLOCMC_CUDA_CHECKED_CALL(cudaGetSymbolAddress((void**)&heap,obj)); ScatterKernelDetail::initKernel<<<1,256>>>(heap, pool, memsize); return heap; } @@ -795,7 +795,7 @@ namespace ScatterKernelDetail{ template < typename T_Obj > static void finalizeHeap(const T_Obj& obj, void* pool){ T_Obj* heap; - SCATTERALLOC_CUDA_CHECKED_CALL(cudaGetSymbolAddress((void**)&heap,obj)); + MALLOCMC_CUDA_CHECKED_CALL(cudaGetSymbolAddress((void**)&heap,obj)); ScatterKernelDetail::finalizeKernel<<<1,256>>>(heap); } @@ -900,7 +900,7 @@ namespace ScatterKernelDetail{ template static unsigned getAvailableSlotsHost(size_t const slotSize, const T_Obj& obj){ T_Obj* heap; - SCATTERALLOC_CUDA_CHECKED_CALL(cudaGetSymbolAddress((void**)&heap,obj)); + MALLOCMC_CUDA_CHECKED_CALL(cudaGetSymbolAddress((void**)&heap,obj)); unsigned h_slots = 0; unsigned* d_slots; cudaMalloc((void**) &d_slots, sizeof(unsigned)); @@ -968,4 +968,4 @@ namespace ScatterKernelDetail{ }; } //namespace CreationPolicies -} //namespace PolicyMalloc +} //namespace mallocMC diff --git a/src/include/scatteralloc/distributionPolicies/Noop.hpp b/src/include/mallocMC/distributionPolicies/Noop.hpp similarity index 92% rename from src/include/scatteralloc/distributionPolicies/Noop.hpp rename to src/include/mallocMC/distributionPolicies/Noop.hpp index d83122bb..608a5f2f 100644 --- a/src/include/scatteralloc/distributionPolicies/Noop.hpp +++ b/src/include/mallocMC/distributionPolicies/Noop.hpp @@ -1,5 +1,5 @@ /* - ScatterAlloc: Massively Parallel Dynamic Memory Allocation for the GPU. + mallocMC: Memory Allocator for Many Core Architectures. Copyright 2014 Institute of Radiation Physics, Helmholtz-Zentrum Dresden - Rossendorf @@ -28,7 +28,7 @@ #pragma once -namespace PolicyMalloc{ +namespace mallocMC{ namespace DistributionPolicies{ /** @@ -40,4 +40,4 @@ namespace DistributionPolicies{ class Noop; } //namespace DistributionPolicies -} //namespace PolicyMalloc +} //namespace mallocMC diff --git a/src/include/scatteralloc/distributionPolicies/Noop_impl.hpp b/src/include/mallocMC/distributionPolicies/Noop_impl.hpp similarity index 88% rename from src/include/scatteralloc/distributionPolicies/Noop_impl.hpp rename to src/include/mallocMC/distributionPolicies/Noop_impl.hpp index 6ec4681c..6d55c2b9 100644 --- a/src/include/scatteralloc/distributionPolicies/Noop_impl.hpp +++ b/src/include/mallocMC/distributionPolicies/Noop_impl.hpp @@ -1,5 +1,5 @@ /* - ScatterAlloc: Massively Parallel Dynamic Memory Allocation for the GPU. + mallocMC: Memory Allocator for Many Core Architectures. Copyright 2014 Institute of Radiation Physics, Helmholtz-Zentrum Dresden - Rossendorf @@ -31,9 +31,9 @@ #include #include "Noop.hpp" -#include "../policy_malloc_prefixes.hpp" +#include "../mallocMC_prefixes.hpp" -namespace PolicyMalloc{ +namespace mallocMC{ namespace DistributionPolicies{ class Noop @@ -42,12 +42,12 @@ namespace DistributionPolicies{ public: - PMMA_ACCELERATOR + MAMC_ACCELERATOR uint32 collect(uint32 bytes){ return bytes; } - PMMA_ACCELERATOR + MAMC_ACCELERATOR void* distribute(void* allocatedMem){ return allocatedMem; } @@ -59,4 +59,4 @@ namespace DistributionPolicies{ }; } //namespace DistributionPolicies -} //namespace PolicyMalloc +} //namespace mallocMC diff --git a/src/include/scatteralloc/distributionPolicies/XMallocSIMD.hpp b/src/include/mallocMC/distributionPolicies/XMallocSIMD.hpp similarity index 95% rename from src/include/scatteralloc/distributionPolicies/XMallocSIMD.hpp rename to src/include/mallocMC/distributionPolicies/XMallocSIMD.hpp index 12773e5e..c80b785a 100644 --- a/src/include/scatteralloc/distributionPolicies/XMallocSIMD.hpp +++ b/src/include/mallocMC/distributionPolicies/XMallocSIMD.hpp @@ -1,5 +1,5 @@ /* - ScatterAlloc: Massively Parallel Dynamic Memory Allocation for the GPU. + mallocMC: Memory Allocator for Many Core Architectures. http://www.icg.tugraz.at/project/mvp Copyright (C) 2012 Institute for Computer Graphics and Vision, @@ -35,7 +35,7 @@ #include -namespace PolicyMalloc{ +namespace mallocMC{ namespace DistributionPolicies{ namespace XMallocSIMDConf{ @@ -68,4 +68,4 @@ namespace DistributionPolicies{ } //namespace DistributionPolicies -} //namespace PolicyMalloc +} //namespace mallocMC diff --git a/src/include/scatteralloc/distributionPolicies/XMallocSIMD_impl.hpp b/src/include/mallocMC/distributionPolicies/XMallocSIMD_impl.hpp similarity index 88% rename from src/include/scatteralloc/distributionPolicies/XMallocSIMD_impl.hpp rename to src/include/mallocMC/distributionPolicies/XMallocSIMD_impl.hpp index 77f4450c..2c96cd93 100644 --- a/src/include/scatteralloc/distributionPolicies/XMallocSIMD_impl.hpp +++ b/src/include/mallocMC/distributionPolicies/XMallocSIMD_impl.hpp @@ -1,5 +1,5 @@ /* - ScatterAlloc: Massively Parallel Dynamic Memory Allocation for the GPU. + mallocMC: Memory Allocator for Many Core Architectures. http://www.icg.tugraz.at/project/mvp Copyright (C) 2012 Institute for Computer Graphics and Vision, @@ -39,11 +39,11 @@ #include #include -#include "../policy_malloc_utils.hpp" -#include "../policy_malloc_prefixes.hpp" +#include "../mallocMC_utils.hpp" +#include "../mallocMC_prefixes.hpp" #include "XMallocSIMD.hpp" -namespace PolicyMalloc{ +namespace mallocMC{ namespace DistributionPolicies{ template @@ -66,14 +66,14 @@ namespace DistributionPolicies{ * shipped default-parameters (in the inherited struct) have lowest precedence. * They will be overridden by a given configuration struct. However, even the * given configuration struct can be overridden by compile-time command line - * parameters (e.g. -D POLICYMALLOC_DP_XMALLOCSIMD_PAGESIZE 1024) + * parameters (e.g. -D MALLOCMC_DP_XMALLOCSIMD_PAGESIZE 1024) * * default-struct < template-struct < command-line parameter */ -#ifndef POLICYMALLOC_DP_XMALLOCSIMD_PAGESIZE -#define POLICYMALLOC_DP_XMALLOCSIMD_PAGESIZE Properties::pagesize::value +#ifndef MALLOCMC_DP_XMALLOCSIMD_PAGESIZE +#define MALLOCMC_DP_XMALLOCSIMD_PAGESIZE Properties::pagesize::value #endif - static const uint32 pagesize = POLICYMALLOC_DP_XMALLOCSIMD_PAGESIZE; + static const uint32 pagesize = MALLOCMC_DP_XMALLOCSIMD_PAGESIZE; //all the properties must be unsigned integers > 0 BOOST_STATIC_ASSERT(!std::numeric_limits::is_signed); @@ -82,11 +82,11 @@ namespace DistributionPolicies{ public: static const uint32 _pagesize = pagesize; - PMMA_ACCELERATOR + MAMC_ACCELERATOR uint32 collect(uint32 bytes){ can_use_coalescing = false; - warpid = PolicyMalloc::warpid(); + warpid = mallocMC::warpid(); myoffset = 0; threadcount = 0; @@ -113,7 +113,7 @@ namespace DistributionPolicies{ } - PMMA_ACCELERATOR + MAMC_ACCELERATOR void* distribute(void* allocatedMem){ __shared__ char* warp_res[32]; @@ -137,7 +137,7 @@ namespace DistributionPolicies{ return myres; } - PMMA_HOST + MAMC_HOST static std::string classname(){ std::stringstream ss; ss << "XMallocSIMD[" << pagesize << "]"; @@ -148,4 +148,4 @@ namespace DistributionPolicies{ } //namespace DistributionPolicies -} //namespace PolicyMalloc +} //namespace mallocMC diff --git a/src/include/scatteralloc/scatteralloc.hpp b/src/include/mallocMC/mallocMC.hpp similarity index 92% rename from src/include/scatteralloc/scatteralloc.hpp rename to src/include/mallocMC/mallocMC.hpp index f33526be..98b7f04d 100644 --- a/src/include/scatteralloc/scatteralloc.hpp +++ b/src/include/mallocMC/mallocMC.hpp @@ -1,5 +1,5 @@ /* - ScatterAlloc: Massively Parallel Dynamic Memory Allocation for the GPU. + mallocMC: Memory Allocator for Many Core Architectures. http://www.icg.tugraz.at/project/mvp https://www.hzdr.de/crp @@ -40,8 +40,8 @@ #include "version.hpp" // core functionality -#include "policy_malloc_overwrites.hpp" -#include "policy_malloc_hostclass.hpp" +#include "mallocMC_overwrites.hpp" +#include "mallocMC_hostclass.hpp" // all the policies #include "CreationPolicies.hpp" diff --git a/src/include/scatteralloc/policy_malloc_config.hpp b/src/include/mallocMC/mallocMC_config.hpp similarity index 62% rename from src/include/scatteralloc/policy_malloc_config.hpp rename to src/include/mallocMC/mallocMC_config.hpp index 0c3e0bd7..40d40e9b 100644 --- a/src/include/scatteralloc/policy_malloc_config.hpp +++ b/src/include/mallocMC/mallocMC_config.hpp @@ -1,5 +1,5 @@ /* - ScatterAlloc: Massively Parallel Dynamic Memory Allocation for the GPU. + mallocMC: Memory Allocator for Many Core Architectures. https://www.hzdr.de/crp Copyright 2014 Institute of Radiation Physics, @@ -31,16 +31,16 @@ #include #include -// basic files for PolicyMalloc -#include "src/include/scatteralloc/policy_malloc_overwrites.hpp" -#include "src/include/scatteralloc/policy_malloc_hostclass.hpp" +// basic files for mallocMC +#include "src/include/mallocMC/mallocMC_overwrites.hpp" +#include "src/include/mallocMC/mallocMC_hostclass.hpp" -// Load all available policies for PolicyMalloc -#include "src/include/scatteralloc/CreationPolicies.hpp" -#include "src/include/scatteralloc/DistributionPolicies.hpp" -#include "src/include/scatteralloc/OOMPolicies.hpp" -#include "src/include/scatteralloc/ReservePoolPolicies.hpp" -#include "src/include/scatteralloc/AlignmentPolicies.hpp" +// Load all available policies for mallocMC +#include "src/include/mallocMC/CreationPolicies.hpp" +#include "src/include/mallocMC/DistributionPolicies.hpp" +#include "src/include/mallocMC/OOMPolicies.hpp" +#include "src/include/mallocMC/ReservePoolPolicies.hpp" +#include "src/include/mallocMC/AlignmentPolicies.hpp" // configurate the CreationPolicy "Scatter" @@ -71,25 +71,25 @@ struct AlignmentConfig{ // Define a new allocator and call it ScatterAllocator // which resembles the behaviour of ScatterAlloc -typedef PolicyMalloc::PolicyAllocator< - PolicyMalloc::CreationPolicies::Scatter, - PolicyMalloc::DistributionPolicies::XMallocSIMD, - PolicyMalloc::OOMPolicies::ReturnNull, - PolicyMalloc::ReservePoolPolicies::SimpleCudaMalloc, - PolicyMalloc::AlignmentPolicies::Shrink +typedef mallocMC::Allocator< + mallocMC::CreationPolicies::Scatter, + mallocMC::DistributionPolicies::XMallocSIMD, + mallocMC::OOMPolicies::ReturnNull, + mallocMC::ReservePoolPolicies::SimpleCudaMalloc, + mallocMC::AlignmentPolicies::Shrink > ScatterAllocator; -typedef PolicyMalloc::PolicyAllocator< - PolicyMalloc::CreationPolicies::OldMalloc, - PolicyMalloc::DistributionPolicies::Noop, - PolicyMalloc::OOMPolicies::ReturnNull, - PolicyMalloc::ReservePoolPolicies::CudaSetLimits, - PolicyMalloc::AlignmentPolicies::Noop +typedef mallocMC::Allocator< + mallocMC::CreationPolicies::OldMalloc, + mallocMC::DistributionPolicies::Noop, + mallocMC::OOMPolicies::ReturnNull, + mallocMC::ReservePoolPolicies::CudaSetLimits, + mallocMC::AlignmentPolicies::Noop > OldAllocator; -// use "ScatterAllocator" as PolicyAllocator -POLICYMALLOC_SET_ALLOCATOR_TYPE(ScatterAllocator) +// use "ScatterAllocator" as Allocator +MALLOCMC_SET_ALLOCATOR_TYPE(ScatterAllocator) -//POLICYMALLOC_SET_ALLOCATOR_TYPE(OldAllocator) +//MALLOCMC_SET_ALLOCATOR_TYPE(OldAllocator) -//POLICYMALLOC_OVERWRITE_MALLOC() +//MALLOCMC_OVERWRITE_MALLOC() diff --git a/src/include/scatteralloc/policy_malloc_constraints.hpp b/src/include/mallocMC/mallocMC_constraints.hpp similarity index 95% rename from src/include/scatteralloc/policy_malloc_constraints.hpp rename to src/include/mallocMC/mallocMC_constraints.hpp index b56d46f1..7aedf4bf 100644 --- a/src/include/scatteralloc/policy_malloc_constraints.hpp +++ b/src/include/mallocMC/mallocMC_constraints.hpp @@ -1,5 +1,5 @@ /* - ScatterAlloc: Massively Parallel Dynamic Memory Allocation for the GPU. + mallocMC: Memory Allocator for Many Core Architectures. https://www.hzdr.de/crp Copyright 2014 Institute of Radiation Physics, @@ -32,7 +32,7 @@ #include "distributionPolicies/XMallocSIMD.hpp" #include -namespace PolicyMalloc{ +namespace mallocMC{ /** The default PolicyCheckers (do always succeed) */ @@ -83,4 +83,4 @@ namespace PolicyMalloc{ Pagesize_must_be_the_same_when_combining_Scatter_and_XMallocSIMD, () ); }; -}//namespace PolicyMalloc +}//namespace mallocMC diff --git a/src/include/scatteralloc/policy_malloc_hostclass.hpp b/src/include/mallocMC/mallocMC_hostclass.hpp similarity index 91% rename from src/include/scatteralloc/policy_malloc_hostclass.hpp rename to src/include/mallocMC/mallocMC_hostclass.hpp index c739f1c9..fb7528e4 100644 --- a/src/include/scatteralloc/policy_malloc_hostclass.hpp +++ b/src/include/mallocMC/mallocMC_hostclass.hpp @@ -1,5 +1,5 @@ /* - ScatterAlloc: Massively Parallel Dynamic Memory Allocation for the GPU. + mallocMC: Memory Allocator for Many Core Architectures. https://www.hzdr.de/crp Copyright 2014 Institute of Radiation Physics, @@ -28,9 +28,9 @@ #pragma once -#include "policy_malloc_utils.hpp" -#include "policy_malloc_constraints.hpp" -#include "policy_malloc_prefixes.hpp" +#include "mallocMC_utils.hpp" +#include "mallocMC_constraints.hpp" +#include "mallocMC_prefixes.hpp" #include #include @@ -40,10 +40,10 @@ #include -namespace PolicyMalloc{ +namespace mallocMC{ /** - * @brief Defines Traits for certain PolicyAllocators + * @brief Defines Traits for certain Allocators * * This trait class provides information about the capabilities of the * allocator. @@ -79,7 +79,7 @@ namespace PolicyMalloc{ typename T_ReservePoolPolicy, typename T_AlignmentPolicy > - struct PolicyAllocator : + struct Allocator : public T_CreationPolicy, public T_OOMPolicy, public T_ReservePoolPolicy, @@ -102,10 +102,10 @@ namespace PolicyMalloc{ public: - typedef PolicyAllocator MyType; - PMMA_ACCELERATOR + MAMC_ACCELERATOR void* alloc(size_t bytes){ DistributionPolicy distributionPolicy; @@ -122,12 +122,12 @@ namespace PolicyMalloc{ // } } - PMMA_ACCELERATOR + MAMC_ACCELERATOR void dealloc(void* p){ CreationPolicy::destroy(p); } - PMMA_HOST + MAMC_HOST void* initHeap(size_t size){ pool = ReservePoolPolicy::setMemPool(size); boost::tie(pool,size) = AlignmentPolicy::alignPool(pool,size); @@ -149,13 +149,13 @@ namespace PolicyMalloc{ return h; } - PMMA_HOST + MAMC_HOST void finalizeHeap(){ CreationPolicy::finalizeHeap(*this,pool); ReservePoolPolicy::resetMemPool(pool); } - PMMA_HOST + MAMC_HOST static std::string info(std::string linebreak = " "){ std::stringstream ss; ss << "CreationPolicy: " << CreationPolicy::classname() << linebreak; @@ -168,18 +168,18 @@ namespace PolicyMalloc{ // polymorphism over the availability of getAvailableSlots - PMMA_HOST PMMA_ACCELERATOR + MAMC_HOST MAMC_ACCELERATOR unsigned getAvailableSlots(size_t slotSize){ return getAvailSlotsPoly(slotSize, boost::mpl::bool_()); } private: - PMMA_HOST PMMA_ACCELERATOR + MAMC_HOST MAMC_ACCELERATOR unsigned getAvailSlotsPoly(size_t slotSize, boost::mpl::bool_){ return 0; } - PMMA_HOST PMMA_ACCELERATOR + MAMC_HOST MAMC_ACCELERATOR unsigned getAvailSlotsPoly(size_t slotSize, boost::mpl::bool_){ slotSize = AlignmentPolicy::applyPadding(slotSize); #ifdef __CUDA_ARCH__ @@ -191,4 +191,4 @@ namespace PolicyMalloc{ }; -} //namespace PolicyMalloc +} //namespace mallocMC diff --git a/src/include/scatteralloc/policy_malloc_overwrites.hpp b/src/include/mallocMC/mallocMC_overwrites.hpp similarity index 60% rename from src/include/scatteralloc/policy_malloc_overwrites.hpp rename to src/include/mallocMC/mallocMC_overwrites.hpp index 712e6f31..97c2c52a 100644 --- a/src/include/scatteralloc/policy_malloc_overwrites.hpp +++ b/src/include/mallocMC/mallocMC_overwrites.hpp @@ -1,5 +1,5 @@ /* - ScatterAlloc: Massively Parallel Dynamic Memory Allocation for the GPU. + mallocMC: Memory Allocator for Many Core Architectures. http://www.icg.tugraz.at/project/mvp https://www.hzdr.de/crp @@ -32,76 +32,76 @@ #pragma once -#include "policy_malloc_prefixes.hpp" +#include "mallocMC_prefixes.hpp" -/** Creates a global object of a policyBased memory allocator +/** Creates a global object of a many core memory allocator * * Will create a global object of the supplied type and use it to generate * global functions that use this type internally. This should be done after - * defining a new policy based memory allocator with a typedef. + * defining a new many core memory allocator with a typedef. */ -#define POLICYMALLOC_GLOBAL_FUNCTIONS(POLICYMALLOC_USER_DEFINED_TYPENAME) \ -namespace PolicyMalloc{ \ - typedef POLICYMALLOC_USER_DEFINED_TYPENAME PolicyMallocType; \ +#define MALLOCMC_GLOBAL_FUNCTIONS(MALLOCMC_USER_DEFINED_TYPENAME) \ +namespace mallocMC{ \ + typedef MALLOCMC_USER_DEFINED_TYPENAME mallocMCType; \ \ -PMMA_ACCELERATOR PolicyMallocType policyMallocGlobalObject; \ +MAMC_ACCELERATOR mallocMCType mallocMCGlobalObject; \ \ -PMMA_HOST void* initHeap( \ +MAMC_HOST void* initHeap( \ size_t heapsize = 8U*1024U*1024U, \ - PolicyMallocType &p = policyMallocGlobalObject \ + mallocMCType &p = mallocMCGlobalObject \ ) \ { \ return p.initHeap(heapsize); \ } \ -PMMA_HOST void finalizeHeap( \ - PolicyMallocType &p = policyMallocGlobalObject \ +MAMC_HOST void finalizeHeap( \ + mallocMCType &p = mallocMCGlobalObject \ ) \ { \ p.finalizeHeap(); \ } \ -} /* end namespace PolicyMalloc */ +} /* end namespace mallocMC */ /** Provides the easily accessible functions getAvaliableSlots * - * Will use the global object defined by POLICYMALLOC_SET_ALLOCATOR_TYPE and + * Will use the global object defined by MALLOCMC_SET_ALLOCATOR_TYPE and * use it to generate global functions that use this type internally. */ -#define POLICYMALLOC_AVAILABLESLOTS() \ -namespace PolicyMalloc{ \ -PMMA_HOST PMMA_ACCELERATOR \ +#define MALLOCMC_AVAILABLESLOTS() \ +namespace mallocMC{ \ +MAMC_HOST MAMC_ACCELERATOR \ unsigned getAvailableSlots( \ size_t slotSize, \ - PolicyMallocType &p = policyMallocGlobalObject){ \ + mallocMCType &p = mallocMCGlobalObject){ \ return p.getAvailableSlots(slotSize); \ } \ -PMMA_HOST PMMA_ACCELERATOR \ +MAMC_HOST MAMC_ACCELERATOR \ bool providesAvailableSlots(){ \ - return Traits::providesAvailableSlots; \ + return Traits::providesAvailableSlots; \ } \ -} /* end namespace PolicyMalloc */ +} /* end namespace mallocMC */ -/** Create the functions pbMalloc() and pbFree() inside a namespace +/** Create the functions mallocMC() and mcfree() inside a namespace * * This allows to use a function without bothering with name-clashes when * including a namespace in the global scope. It will call the namespaced * version of malloc() inside. */ -#define POLICYMALLOC_PBMALLOC() \ -namespace PolicyMalloc{ \ -PMMA_ACCELERATOR \ -void* pbMalloc(size_t t) __THROW \ +#define MALLOCMC_MALLOCMC() \ +namespace mallocMC{ \ +MAMC_ACCELERATOR \ +void* mallocMC(size_t t) __THROW \ { \ - return PolicyMalloc::malloc(t); \ + return mallocMC::malloc(t); \ } \ -PMMA_ACCELERATOR \ -void pbFree(void* p) __THROW \ +MAMC_ACCELERATOR \ +void mcfree(void* p) __THROW \ { \ - PolicyMalloc::free(p); \ + mallocMC::free(p); \ } \ -} /* end namespace PolicyMalloc */ +} /* end namespace mallocMC */ /** Create the functions malloc() and free() inside a namespace @@ -110,19 +110,19 @@ void pbFree(void* p) __THROW \ * "malloc" or "free". This is useful when using a policy that contains a call * to the original device-side malloc() from CUDA. */ -#define POLICYMALLOC_MALLOC() \ -namespace PolicyMalloc{ \ -PMMA_ACCELERATOR \ +#define MALLOCMC_MALLOC() \ +namespace mallocMC{ \ +MAMC_ACCELERATOR \ void* malloc(size_t t) __THROW \ { \ - return PolicyMalloc::policyMallocGlobalObject.alloc(t); \ + return mallocMC::mallocMCGlobalObject.alloc(t); \ } \ -PMMA_ACCELERATOR \ +MAMC_ACCELERATOR \ void free(void* p) __THROW \ { \ - PolicyMalloc::policyMallocGlobalObject.dealloc(p); \ + mallocMC::mallocMCGlobalObject.dealloc(p); \ } \ -} /* end namespace PolicyMalloc */ +} /* end namespace mallocMC */ @@ -137,24 +137,24 @@ void free(void* p) __THROW \ */ #ifdef __CUDACC__ #if __CUDA_ARCH__ >= 200 -#define POLICYMALLOC_OVERWRITE_NEW() \ -PMMA_ACCELERATOR \ -void* operator new(size_t t, PolicyMalloc::PolicyMallocType &h) \ +#define MALLOCMC_OVERWRITE_NEW() \ +MAMC_ACCELERATOR \ +void* operator new(size_t t, mallocMC::mallocMCType &h) \ { \ return h.alloc(t); \ } \ -PMMA_ACCELERATOR \ -void* operator new[](size_t t, PolicyMalloc::PolicyMallocType &h) \ +MAMC_ACCELERATOR \ +void* operator new[](size_t t, mallocMC::mallocMCType &h) \ { \ return h.alloc(t); \ } \ -PMMA_ACCELERATOR \ -void operator delete(void* p, PolicyMalloc::PolicyMallocType &h) \ +MAMC_ACCELERATOR \ +void operator delete(void* p, mallocMC::mallocMCType &h) \ { \ h.dealloc(p); \ } \ -PMMA_ACCELERATOR \ -void operator delete[](void* p, PolicyMalloc::PolicyMallocType &h) \ +MAMC_ACCELERATOR \ +void operator delete[](void* p, mallocMC::mallocMCType &h) \ { \ h.dealloc(p); \ } @@ -171,16 +171,16 @@ void operator delete[](void* p, PolicyMalloc::PolicyMallocType &h) \ */ #ifdef __CUDACC__ #if __CUDA_ARCH__ >= 200 -#define POLICYMALLOC_OVERWRITE_MALLOC() \ -PMMA_ACCELERATOR \ +#define MALLOCMC_OVERWRITE_MALLOC() \ +MAMC_ACCELERATOR \ void* malloc(size_t t) __THROW \ { \ - return PolicyMalloc::malloc(t); \ + return mallocMC::malloc(t); \ } \ -PMMA_ACCELERATOR \ +MAMC_ACCELERATOR \ void free(void* p) __THROW \ { \ - PolicyMalloc::free(p); \ + mallocMC::free(p); \ } #endif #endif @@ -190,20 +190,20 @@ void free(void* p) __THROW \ /* if the defines do not exist (wrong CUDA version etc), * create at least empty defines */ -#ifndef POLICYMALLOC_PBMALLOC -#define POLICYMALLOC_PBMALLOC() +#ifndef MALLOCMC_MALLOCMC +#define MALLOCMC_MALLOCMC() #endif -#ifndef POLICYMALLOC_MALLOC -#define POLICYMALLOC_MALLOC() +#ifndef MALLOCMC_MALLOC +#define MALLOCMC_MALLOC() #endif -#ifndef POLICYMALLOC_OVERWRITE_NEW -#define POLICYMALLOC_OVERWRITE_NEW() +#ifndef MALLOCMC_OVERWRITE_NEW +#define MALLOCMC_OVERWRITE_NEW() #endif -#ifndef POLICYMALLOC_OVERWRITE_MALLOC -#define POLICYMALLOC_OVERWRITE_MALLOC() +#ifndef MALLOCMC_OVERWRITE_MALLOC +#define MALLOCMC_OVERWRITE_MALLOC() #endif @@ -214,11 +214,11 @@ void free(void* p) __THROW \ * will create a global object of that allocator and the necessary functions to * initialize and allocate memory. */ -#define POLICYMALLOC_SET_ALLOCATOR_TYPE(POLICYMALLOC_USER_DEFINED_TYPE) \ -POLICYMALLOC_GLOBAL_FUNCTIONS(POLICYMALLOC_USER_DEFINED_TYPE) \ -POLICYMALLOC_MALLOC() \ -POLICYMALLOC_PBMALLOC() \ -POLICYMALLOC_AVAILABLESLOTS() +#define MALLOCMC_SET_ALLOCATOR_TYPE(MALLOCMC_USER_DEFINED_TYPE) \ +MALLOCMC_GLOBAL_FUNCTIONS(MALLOCMC_USER_DEFINED_TYPE) \ +MALLOCMC_MALLOC() \ +MALLOCMC_MALLOCMC() \ +MALLOCMC_AVAILABLESLOTS() -//POLICYMALLOC_OVERWRITE_NEW() +//MALLOCMC_OVERWRITE_NEW() diff --git a/src/include/scatteralloc/policy_malloc_prefixes.hpp b/src/include/mallocMC/mallocMC_prefixes.hpp similarity index 90% rename from src/include/scatteralloc/policy_malloc_prefixes.hpp rename to src/include/mallocMC/mallocMC_prefixes.hpp index 9eb7c449..1199e3e1 100644 --- a/src/include/scatteralloc/policy_malloc_prefixes.hpp +++ b/src/include/mallocMC/mallocMC_prefixes.hpp @@ -1,5 +1,5 @@ /* - ScatterAlloc: Massively Parallel Dynamic Memory Allocation for the GPU. + mallocMC: Memory Allocator for Many Core Architectures. https://www.hzdr.de/crp Copyright (C) 2014 Institute of Radiation Physics, @@ -28,6 +28,6 @@ #pragma once -#define PMMA_HOST __host__ -#define PMMA_ACCELERATOR __device__ +#define MAMC_HOST __host__ +#define MAMC_ACCELERATOR __device__ diff --git a/src/include/scatteralloc/policy_malloc_utils.hpp b/src/include/mallocMC/mallocMC_utils.hpp similarity index 80% rename from src/include/scatteralloc/policy_malloc_utils.hpp rename to src/include/mallocMC/mallocMC_utils.hpp index 159e72f9..0353952c 100644 --- a/src/include/scatteralloc/policy_malloc_utils.hpp +++ b/src/include/mallocMC/mallocMC_utils.hpp @@ -1,5 +1,5 @@ /* - ScatterAlloc: Massively Parallel Dynamic Memory Allocation for the GPU. + mallocMC: Memory Allocator for Many Core Architectures. http://www.icg.tugraz.at/project/mvp https://www.hzdr.de/crp @@ -41,7 +41,7 @@ #include #include -#include "policy_malloc_prefixes.hpp" +#include "mallocMC_prefixes.hpp" namespace CUDA @@ -94,15 +94,15 @@ namespace CUDA #endif } -#define SCATTERALLOC_CUDA_CHECKED_CALL(call) CUDA::checkError(call, __FILE__, __LINE__) -#define SCATTERALLOC_CUDA_CHECK_ERROR() CUDA::checkError(__FILE__, __LINE__) +#define MALLOCMC_CUDA_CHECKED_CALL(call) CUDA::checkError(call, __FILE__, __LINE__) +#define MALLOCMC_CUDA_CHECK_ERROR() CUDA::checkError(__FILE__, __LINE__) } #define warp_serial \ for (uint32_richtig_huebsch __mask = __ballot(1), \ __num = __popc(__mask), \ - __lanemask = PolicyMalloc::lanemask_lt(), \ + __lanemask = mallocMC::lanemask_lt(), \ __local_id = __popc(__lanemask & __mask), \ __active = 0; \ __active < __num; \ @@ -110,7 +110,7 @@ namespace CUDA if (__active == __local_id) -namespace PolicyMalloc +namespace mallocMC { typedef unsigned int uint32_richtig_huebsch; @@ -127,71 +127,71 @@ namespace PolicyMalloc typedef unsigned long long int type; }; - typedef PolicyMalloc::__PointerEquivalent::type PointerEquivalent; + typedef mallocMC::__PointerEquivalent::type PointerEquivalent; - PMMA_ACCELERATOR inline uint32_richtig_huebsch laneid() + MAMC_ACCELERATOR inline uint32_richtig_huebsch laneid() { uint32_richtig_huebsch mylaneid; asm("mov.u32 %0, %laneid;" : "=r" (mylaneid)); return mylaneid; } - PMMA_ACCELERATOR inline uint32_richtig_huebsch warpid() + MAMC_ACCELERATOR inline uint32_richtig_huebsch warpid() { uint32_richtig_huebsch mywarpid; asm("mov.u32 %0, %warpid;" : "=r" (mywarpid)); return mywarpid; } - PMMA_ACCELERATOR inline uint32_richtig_huebsch nwarpid() + MAMC_ACCELERATOR inline uint32_richtig_huebsch nwarpid() { uint32_richtig_huebsch mynwarpid; asm("mov.u32 %0, %nwarpid;" : "=r" (mynwarpid)); return mynwarpid; } - PMMA_ACCELERATOR inline uint32_richtig_huebsch smid() + MAMC_ACCELERATOR inline uint32_richtig_huebsch smid() { uint32_richtig_huebsch mysmid; asm("mov.u32 %0, %smid;" : "=r" (mysmid)); return mysmid; } - PMMA_ACCELERATOR inline uint32_richtig_huebsch nsmid() + MAMC_ACCELERATOR inline uint32_richtig_huebsch nsmid() { uint32_richtig_huebsch mynsmid; asm("mov.u32 %0, %nsmid;" : "=r" (mynsmid)); return mynsmid; } - PMMA_ACCELERATOR inline uint32_richtig_huebsch lanemask() + MAMC_ACCELERATOR inline uint32_richtig_huebsch lanemask() { uint32_richtig_huebsch lanemask; asm("mov.u32 %0, %lanemask_eq;" : "=r" (lanemask)); return lanemask; } - PMMA_ACCELERATOR inline uint32_richtig_huebsch lanemask_le() + MAMC_ACCELERATOR inline uint32_richtig_huebsch lanemask_le() { uint32_richtig_huebsch lanemask; asm("mov.u32 %0, %lanemask_le;" : "=r" (lanemask)); return lanemask; } - PMMA_ACCELERATOR inline uint32_richtig_huebsch lanemask_lt() + MAMC_ACCELERATOR inline uint32_richtig_huebsch lanemask_lt() { uint32_richtig_huebsch lanemask; asm("mov.u32 %0, %lanemask_lt;" : "=r" (lanemask)); return lanemask; } - PMMA_ACCELERATOR inline uint32_richtig_huebsch lanemask_ge() + MAMC_ACCELERATOR inline uint32_richtig_huebsch lanemask_ge() { uint32_richtig_huebsch lanemask; asm("mov.u32 %0, %lanemask_ge;" : "=r" (lanemask)); return lanemask; } - PMMA_ACCELERATOR inline uint32_richtig_huebsch lanemask_gt() + MAMC_ACCELERATOR inline uint32_richtig_huebsch lanemask_gt() { uint32_richtig_huebsch lanemask; asm("mov.u32 %0, %lanemask_gt;" : "=r" (lanemask)); @@ -199,6 +199,6 @@ namespace PolicyMalloc } template - PMMA_HOST PMMA_ACCELERATOR inline T divup(T a, T b) { return (a + b - 1)/b; } + MAMC_HOST MAMC_ACCELERATOR inline T divup(T a, T b) { return (a + b - 1)/b; } } diff --git a/src/include/scatteralloc/oOMPolicies/BadAllocException.hpp b/src/include/mallocMC/oOMPolicies/BadAllocException.hpp similarity index 93% rename from src/include/scatteralloc/oOMPolicies/BadAllocException.hpp rename to src/include/mallocMC/oOMPolicies/BadAllocException.hpp index 85a24d1d..a99be38a 100644 --- a/src/include/scatteralloc/oOMPolicies/BadAllocException.hpp +++ b/src/include/mallocMC/oOMPolicies/BadAllocException.hpp @@ -1,5 +1,5 @@ /* - ScatterAlloc: Massively Parallel Dynamic Memory Allocation for the GPU. + mallocMC: Memory Allocator for Many Core Architectures. Copyright 2014 Institute of Radiation Physics, Helmholtz-Zentrum Dresden - Rossendorf @@ -27,7 +27,7 @@ #pragma once -namespace PolicyMalloc{ +namespace mallocMC{ namespace OOMPolicies{ /** @@ -42,4 +42,4 @@ namespace OOMPolicies{ class BadAllocException; } //namespace OOMPolicies -} //namespace PolicyMalloc +} //namespace mallocMC diff --git a/src/include/scatteralloc/oOMPolicies/BadAllocException_impl.hpp b/src/include/mallocMC/oOMPolicies/BadAllocException_impl.hpp similarity index 90% rename from src/include/scatteralloc/oOMPolicies/BadAllocException_impl.hpp rename to src/include/mallocMC/oOMPolicies/BadAllocException_impl.hpp index 7c8c1225..447c254c 100644 --- a/src/include/scatteralloc/oOMPolicies/BadAllocException_impl.hpp +++ b/src/include/mallocMC/oOMPolicies/BadAllocException_impl.hpp @@ -1,5 +1,5 @@ /* - ScatterAlloc: Massively Parallel Dynamic Memory Allocation for the GPU. + mallocMC: Memory Allocator for Many Core Architectures. Copyright 2014 Institute of Radiation Physics, Helmholtz-Zentrum Dresden - Rossendorf @@ -31,14 +31,14 @@ #include #include "BadAllocException.hpp" -#include "../policy_malloc_prefixes.hpp" +#include "../mallocMC_prefixes.hpp" -namespace PolicyMalloc{ +namespace mallocMC{ namespace OOMPolicies{ struct BadAllocException { - PMMA_ACCELERATOR + MAMC_ACCELERATOR static void* handleOOM(void* mem){ #ifdef __CUDACC__ //#if __CUDA_ARCH__ < 350 @@ -62,4 +62,4 @@ namespace OOMPolicies{ }; } //namespace OOMPolicies -} //namespace PolicyMalloc +} //namespace mallocMC diff --git a/src/include/scatteralloc/oOMPolicies/ReturnNull.hpp b/src/include/mallocMC/oOMPolicies/ReturnNull.hpp similarity index 92% rename from src/include/scatteralloc/oOMPolicies/ReturnNull.hpp rename to src/include/mallocMC/oOMPolicies/ReturnNull.hpp index ee7a3274..2db8568a 100644 --- a/src/include/scatteralloc/oOMPolicies/ReturnNull.hpp +++ b/src/include/mallocMC/oOMPolicies/ReturnNull.hpp @@ -1,5 +1,5 @@ /* - ScatterAlloc: Massively Parallel Dynamic Memory Allocation for the GPU. + mallocMC: Memory Allocator for Many Core Architectures. Copyright 2014 Institute of Radiation Physics, Helmholtz-Zentrum Dresden - Rossendorf @@ -27,7 +27,7 @@ #pragma once -namespace PolicyMalloc{ +namespace mallocMC{ namespace OOMPolicies{ /** @@ -38,4 +38,4 @@ namespace OOMPolicies{ class ReturnNull; } //namespace OOMPolicies -} //namespace PolicyMalloc +} //namespace mallocMC diff --git a/src/include/scatteralloc/oOMPolicies/ReturnNull_impl.hpp b/src/include/mallocMC/oOMPolicies/ReturnNull_impl.hpp similarity index 88% rename from src/include/scatteralloc/oOMPolicies/ReturnNull_impl.hpp rename to src/include/mallocMC/oOMPolicies/ReturnNull_impl.hpp index d36ac5e2..f18bd0e2 100644 --- a/src/include/scatteralloc/oOMPolicies/ReturnNull_impl.hpp +++ b/src/include/mallocMC/oOMPolicies/ReturnNull_impl.hpp @@ -1,5 +1,5 @@ /* - ScatterAlloc: Massively Parallel Dynamic Memory Allocation for the GPU. + mallocMC: Memory Allocator for Many Core Architectures. Copyright 2014 Institute of Radiation Physics, Helmholtz-Zentrum Dresden - Rossendorf @@ -30,15 +30,15 @@ #include #include "ReturnNull.hpp" -#include "../policy_malloc_prefixes.hpp" +#include "../mallocMC_prefixes.hpp" -namespace PolicyMalloc{ +namespace mallocMC{ namespace OOMPolicies{ class ReturnNull { public: - PMMA_ACCELERATOR + MAMC_ACCELERATOR static void* handleOOM(void* mem){ return NULL; } @@ -49,4 +49,4 @@ namespace OOMPolicies{ }; } //namespace OOMPolicies -} //namespace PolicyMalloc +} //namespace mallocMC diff --git a/src/include/scatteralloc/reservePoolPolicies/CudaSetLimits.hpp b/src/include/mallocMC/reservePoolPolicies/CudaSetLimits.hpp similarity index 93% rename from src/include/scatteralloc/reservePoolPolicies/CudaSetLimits.hpp rename to src/include/mallocMC/reservePoolPolicies/CudaSetLimits.hpp index db6d18af..cfbceea2 100644 --- a/src/include/scatteralloc/reservePoolPolicies/CudaSetLimits.hpp +++ b/src/include/mallocMC/reservePoolPolicies/CudaSetLimits.hpp @@ -1,5 +1,5 @@ /* - ScatterAlloc: Massively Parallel Dynamic Memory Allocation for the GPU. + mallocMC: Memory Allocator for Many Core Architectures. Copyright 2014 Institute of Radiation Physics, Helmholtz-Zentrum Dresden - Rossendorf @@ -27,7 +27,7 @@ #pragma once -namespace PolicyMalloc{ +namespace mallocMC{ namespace ReservePoolPolicies{ /** @@ -41,4 +41,4 @@ namespace ReservePoolPolicies{ struct CudaSetLimits; } //namespace ReservePoolPolicies -} //namespace PolicyMalloc +} //namespace mallocMC diff --git a/src/include/scatteralloc/reservePoolPolicies/CudaSetLimits_impl.hpp b/src/include/mallocMC/reservePoolPolicies/CudaSetLimits_impl.hpp similarity index 93% rename from src/include/scatteralloc/reservePoolPolicies/CudaSetLimits_impl.hpp rename to src/include/mallocMC/reservePoolPolicies/CudaSetLimits_impl.hpp index fd87d5b8..eb4ddfa2 100644 --- a/src/include/scatteralloc/reservePoolPolicies/CudaSetLimits_impl.hpp +++ b/src/include/mallocMC/reservePoolPolicies/CudaSetLimits_impl.hpp @@ -1,5 +1,5 @@ /* - ScatterAlloc: Massively Parallel Dynamic Memory Allocation for the GPU. + mallocMC: Memory Allocator for Many Core Architectures. Copyright 2014 Institute of Radiation Physics, Helmholtz-Zentrum Dresden - Rossendorf @@ -32,7 +32,7 @@ #include "CudaSetLimits.hpp" -namespace PolicyMalloc{ +namespace mallocMC{ namespace ReservePoolPolicies{ struct CudaSetLimits{ @@ -52,4 +52,4 @@ namespace ReservePoolPolicies{ }; } //namespace ReservePoolPolicies -} //namespace PolicyMalloc +} //namespace mallocMC diff --git a/src/include/scatteralloc/reservePoolPolicies/SimpleCudaMalloc.hpp b/src/include/mallocMC/reservePoolPolicies/SimpleCudaMalloc.hpp similarity index 93% rename from src/include/scatteralloc/reservePoolPolicies/SimpleCudaMalloc.hpp rename to src/include/mallocMC/reservePoolPolicies/SimpleCudaMalloc.hpp index 008ea111..1bb386f7 100644 --- a/src/include/scatteralloc/reservePoolPolicies/SimpleCudaMalloc.hpp +++ b/src/include/mallocMC/reservePoolPolicies/SimpleCudaMalloc.hpp @@ -1,5 +1,5 @@ /* - ScatterAlloc: Massively Parallel Dynamic Memory Allocation for the GPU. + mallocMC: Memory Allocator for Many Core Architectures. Copyright 2014 Institute of Radiation Physics, Helmholtz-Zentrum Dresden - Rossendorf @@ -27,7 +27,7 @@ #pragma once -namespace PolicyMalloc{ +namespace mallocMC{ namespace ReservePoolPolicies{ /** @@ -41,4 +41,4 @@ namespace ReservePoolPolicies{ struct SimpleCudaMalloc; } //namespace ReservePoolPolicies -} //namespace PolicyMalloc +} //namespace mallocMC diff --git a/src/include/scatteralloc/reservePoolPolicies/SimpleCudaMalloc_impl.hpp b/src/include/mallocMC/reservePoolPolicies/SimpleCudaMalloc_impl.hpp similarity index 85% rename from src/include/scatteralloc/reservePoolPolicies/SimpleCudaMalloc_impl.hpp rename to src/include/mallocMC/reservePoolPolicies/SimpleCudaMalloc_impl.hpp index bd7f0ac5..da13f4a4 100644 --- a/src/include/scatteralloc/reservePoolPolicies/SimpleCudaMalloc_impl.hpp +++ b/src/include/mallocMC/reservePoolPolicies/SimpleCudaMalloc_impl.hpp @@ -1,5 +1,5 @@ /* - ScatterAlloc: Massively Parallel Dynamic Memory Allocation for the GPU. + mallocMC: Memory Allocator for Many Core Architectures. Copyright 2014 Institute of Radiation Physics, Helmholtz-Zentrum Dresden - Rossendorf @@ -29,22 +29,22 @@ #include -#include "../policy_malloc_utils.hpp" +#include "../mallocMC_utils.hpp" #include "SimpleCudaMalloc.hpp" -namespace PolicyMalloc{ +namespace mallocMC{ namespace ReservePoolPolicies{ struct SimpleCudaMalloc{ static void* setMemPool(size_t memsize){ void* pool; - SCATTERALLOC_CUDA_CHECKED_CALL(cudaMalloc(&pool, memsize)); + MALLOCMC_CUDA_CHECKED_CALL(cudaMalloc(&pool, memsize)); return pool; } static void resetMemPool(void* p){ - SCATTERALLOC_CUDA_CHECKED_CALL(cudaFree(p)); + MALLOCMC_CUDA_CHECKED_CALL(cudaFree(p)); } static std::string classname(){ @@ -54,4 +54,4 @@ namespace ReservePoolPolicies{ }; } //namespace ReservePoolPolicies -} //namespace PolicyMalloc +} //namespace mallocMC diff --git a/src/include/scatteralloc/version.hpp b/src/include/mallocMC/version.hpp similarity index 96% rename from src/include/scatteralloc/version.hpp rename to src/include/mallocMC/version.hpp index 217a8025..750aa46f 100644 --- a/src/include/scatteralloc/version.hpp +++ b/src/include/mallocMC/version.hpp @@ -1,5 +1,5 @@ /* - ScatterAlloc: Massively Parallel Dynamic Memory Allocation for the GPU. + mallocMC: Memory Allocator for Many Core Architectures. http://www.icg.tugraz.at/project/mvp https://www.hzdr.de/crp diff --git a/tests/verify_heap.cu b/tests/verify_heap.cu index 0579c480..9101a7c9 100644 --- a/tests/verify_heap.cu +++ b/tests/verify_heap.cu @@ -1,5 +1,5 @@ /* - ScatterAlloc: Massively Parallel Dynamic Memory Allocation for the GPU. + mallocMC: Memory Allocator for Many Core Architectures. https://www.hzdr.de/crp Copyright 2014 Institute of Radiation Physics, @@ -47,12 +47,12 @@ #include //include the Heap with the arguments given in the config -#include "src/include/scatteralloc/policy_malloc_utils.hpp" +#include "src/include/mallocMC/mallocMC_utils.hpp" #include "verify_heap_config.hpp" //use ScatterAllocator to replace malloc/free -POLICYMALLOC_SET_ALLOCATOR_TYPE(ScatterAllocator) -POLICYMALLOC_OVERWRITE_MALLOC() +MALLOCMC_SET_ALLOCATOR_TYPE(ScatterAllocator) +MALLOCMC_OVERWRITE_MALLOC() // global variable for verbosity, might change due to user input '--verbose' bool verbose = false; @@ -331,7 +331,7 @@ __global__ void allocAll( unsigned long long sum=0; while(true){ - allocElem_t* p = (allocElem_t*) PolicyMalloc::malloc(sizeof(allocElem_t) * ELEMS_PER_SLOT); + allocElem_t* p = (allocElem_t*) mallocMC::malloc(sizeof(allocElem_t) * ELEMS_PER_SLOT); if(p == NULL) break; size_t pos = atomicAdd(counter,1); @@ -364,7 +364,7 @@ __global__ void deallocAll( while(true){ size_t pos = atomicAdd(counter,1); if(pos >= nSlots) break; - PolicyMalloc::free(data[pos]); + mallocMC::free(data[pos]); } } @@ -386,7 +386,7 @@ __global__ void damageElement(allocElem_t** data){ /** * wrapper function to allocate memory on device * - * allocates memory with scatterAlloc. Returns the number of + * allocates memory with mallocMC. Returns the number of * created elements as well as the sum of these elements * * @param d_testData the datastructure which will hold @@ -411,15 +411,15 @@ void allocate( unsigned long long *d_sum; unsigned long long *d_nSlots; - SCATTERALLOC_CUDA_CHECKED_CALL(cudaMalloc((void**) &d_sum,sizeof(unsigned long long))); - SCATTERALLOC_CUDA_CHECKED_CALL(cudaMalloc((void**) &d_nSlots, sizeof(unsigned long long))); - SCATTERALLOC_CUDA_CHECKED_CALL(cudaMemcpy(d_sum,&zero,sizeof(unsigned long long),cudaMemcpyHostToDevice)); - SCATTERALLOC_CUDA_CHECKED_CALL(cudaMemcpy(d_nSlots,&zero,sizeof(unsigned long long),cudaMemcpyHostToDevice)); + MALLOCMC_CUDA_CHECKED_CALL(cudaMalloc((void**) &d_sum,sizeof(unsigned long long))); + MALLOCMC_CUDA_CHECKED_CALL(cudaMalloc((void**) &d_nSlots, sizeof(unsigned long long))); + MALLOCMC_CUDA_CHECKED_CALL(cudaMemcpy(d_sum,&zero,sizeof(unsigned long long),cudaMemcpyHostToDevice)); + MALLOCMC_CUDA_CHECKED_CALL(cudaMemcpy(d_nSlots,&zero,sizeof(unsigned long long),cudaMemcpyHostToDevice)); CUDA_CHECK_KERNEL_SYNC(allocAll<<>>(d_testData,d_nSlots,d_sum)); - SCATTERALLOC_CUDA_CHECKED_CALL(cudaMemcpy(h_sum,d_sum,sizeof(unsigned long long),cudaMemcpyDeviceToHost)); - SCATTERALLOC_CUDA_CHECKED_CALL(cudaMemcpy(h_nSlots,d_nSlots,sizeof(unsigned long long),cudaMemcpyDeviceToHost)); + MALLOCMC_CUDA_CHECKED_CALL(cudaMemcpy(h_sum,d_sum,sizeof(unsigned long long),cudaMemcpyDeviceToHost)); + MALLOCMC_CUDA_CHECKED_CALL(cudaMemcpy(h_nSlots,d_nSlots,sizeof(unsigned long long),cudaMemcpyDeviceToHost)); cudaFree(d_sum); cudaFree(d_nSlots); dout() << "done" << std::endl; @@ -454,12 +454,12 @@ bool verify( unsigned long long *d_sum; unsigned long long *d_counter; - SCATTERALLOC_CUDA_CHECKED_CALL(cudaMalloc((void**) &d_sum, sizeof(unsigned long long))); - SCATTERALLOC_CUDA_CHECKED_CALL(cudaMalloc((void**) &d_counter, sizeof(unsigned long long))); - SCATTERALLOC_CUDA_CHECKED_CALL(cudaMalloc((void**) &d_correct, sizeof(int))); - SCATTERALLOC_CUDA_CHECKED_CALL(cudaMemcpy(d_sum,&zero,sizeof(unsigned long long),cudaMemcpyHostToDevice)); - SCATTERALLOC_CUDA_CHECKED_CALL(cudaMemcpy(d_counter,&zero,sizeof(unsigned long long),cudaMemcpyHostToDevice)); - SCATTERALLOC_CUDA_CHECKED_CALL(cudaMemcpy(d_correct,&h_correct,sizeof(int),cudaMemcpyHostToDevice)); + MALLOCMC_CUDA_CHECKED_CALL(cudaMalloc((void**) &d_sum, sizeof(unsigned long long))); + MALLOCMC_CUDA_CHECKED_CALL(cudaMalloc((void**) &d_counter, sizeof(unsigned long long))); + MALLOCMC_CUDA_CHECKED_CALL(cudaMalloc((void**) &d_correct, sizeof(int))); + MALLOCMC_CUDA_CHECKED_CALL(cudaMemcpy(d_sum,&zero,sizeof(unsigned long long),cudaMemcpyHostToDevice)); + MALLOCMC_CUDA_CHECKED_CALL(cudaMemcpy(d_counter,&zero,sizeof(unsigned long long),cudaMemcpyHostToDevice)); + MALLOCMC_CUDA_CHECKED_CALL(cudaMemcpy(d_correct,&h_correct,sizeof(int),cudaMemcpyHostToDevice)); // can be replaced by a call to check_content_fast, // if the gaussian sum (see below) is not used and you @@ -471,7 +471,7 @@ bool verify( static_cast(nSlots), d_correct )); - SCATTERALLOC_CUDA_CHECKED_CALL(cudaMemcpy(&h_correct,d_correct,sizeof(int),cudaMemcpyDeviceToHost)); + MALLOCMC_CUDA_CHECKED_CALL(cudaMemcpy(&h_correct,d_correct,sizeof(int),cudaMemcpyDeviceToHost)); // This only works, if the type "allocElem_t" // can hold all the IDs (usually unsigned long long) @@ -479,8 +479,8 @@ bool verify( dout() << "verifying on host..."; unsigned long long h_sum, h_counter; unsigned long long gaussian_sum = (ELEMS_PER_SLOT*nSlots * (ELEMS_PER_SLOT*nSlots-1))/2; - SCATTERALLOC_CUDA_CHECKED_CALL(cudaMemcpy(&h_sum,d_sum,sizeof(unsigned long long),cudaMemcpyDeviceToHost)); - SCATTERALLOC_CUDA_CHECKED_CALL(cudaMemcpy(&h_counter,d_counter,sizeof(unsigned long long),cudaMemcpyDeviceToHost)); + MALLOCMC_CUDA_CHECKED_CALL(cudaMemcpy(&h_sum,d_sum,sizeof(unsigned long long),cudaMemcpyDeviceToHost)); + MALLOCMC_CUDA_CHECKED_CALL(cudaMemcpy(&h_counter,d_counter,sizeof(unsigned long long),cudaMemcpyDeviceToHost)); if(gaussian_sum != h_sum){ dout() << "\nGaussian Sum doesn't match: is " << h_sum; dout() << " (should be " << gaussian_sum << ")" << std::endl; @@ -588,7 +588,7 @@ void print_machine_readable( /** - * Verify the heap allocation of ScatterAlloc + * Verify the heap allocation of mallocMC * * Allocates as much memory as the heap allows. Make sure that allocated * memory actually holds the correct values without corrupting them. Will @@ -636,11 +636,11 @@ bool run_heap_verification( dout() << "maximum of elements: " << maxSlots << std::endl; // initializing the heap - PolicyMalloc::initHeap(heapSize); + mallocMC::initHeap(heapSize); allocElem_t** d_testData; - SCATTERALLOC_CUDA_CHECKED_CALL(cudaMalloc((void**) &d_testData, nPointers*sizeof(allocElem_t*))); + MALLOCMC_CUDA_CHECKED_CALL(cudaMalloc((void**) &d_testData, nPointers*sizeof(allocElem_t*))); - // allocating with scatterAlloc + // allocating with mallocMC unsigned long long usedSlots = 0; unsigned long long sumAllocElems = 0; allocate(d_testData,&usedSlots,&sumAllocElems,blocks,threads); @@ -668,12 +668,12 @@ bool run_heap_verification( // release all memory dout() << "deallocation... "; unsigned long long* d_dealloc_counter; - SCATTERALLOC_CUDA_CHECKED_CALL(cudaMalloc((void**) &d_dealloc_counter, sizeof(unsigned long long))); - SCATTERALLOC_CUDA_CHECKED_CALL(cudaMemcpy(d_dealloc_counter,&zero,sizeof(unsigned long long),cudaMemcpyHostToDevice)); + MALLOCMC_CUDA_CHECKED_CALL(cudaMalloc((void**) &d_dealloc_counter, sizeof(unsigned long long))); + MALLOCMC_CUDA_CHECKED_CALL(cudaMemcpy(d_dealloc_counter,&zero,sizeof(unsigned long long),cudaMemcpyHostToDevice)); CUDA_CHECK_KERNEL_SYNC(deallocAll<<>>(d_testData,d_dealloc_counter,static_cast(usedSlots))); cudaFree(d_dealloc_counter); cudaFree(d_testData); - PolicyMalloc::finalizeHeap(); + mallocMC::finalizeHeap(); dout() << "done "<< std::endl; diff --git a/tests/verify_heap_config.hpp b/tests/verify_heap_config.hpp index 951071c6..4371b123 100644 --- a/tests/verify_heap_config.hpp +++ b/tests/verify_heap_config.hpp @@ -1,5 +1,5 @@ /* - ScatterAlloc: Massively Parallel Dynamic Memory Allocation for the GPU. + mallocMC: Memory Allocator for Many Core Architectures. https://www.hzdr.de/crp Copyright 2014 Institute of Radiation Physics, @@ -31,16 +31,16 @@ #include #include -// basic files for PolicyMalloc -#include "src/include/scatteralloc/policy_malloc_overwrites.hpp" -#include "src/include/scatteralloc/policy_malloc_hostclass.hpp" +// basic files for mallocMC +#include "src/include/mallocMC/mallocMC_overwrites.hpp" +#include "src/include/mallocMC/mallocMC_hostclass.hpp" -// Load all available policies for PolicyMalloc -#include "src/include/scatteralloc/CreationPolicies.hpp" -#include "src/include/scatteralloc/DistributionPolicies.hpp" -#include "src/include/scatteralloc/OOMPolicies.hpp" -#include "src/include/scatteralloc/ReservePoolPolicies.hpp" -#include "src/include/scatteralloc/AlignmentPolicies.hpp" +// Load all available policies for mallocMC +#include "src/include/mallocMC/CreationPolicies.hpp" +#include "src/include/mallocMC/DistributionPolicies.hpp" +#include "src/include/mallocMC/OOMPolicies.hpp" +#include "src/include/mallocMC/ReservePoolPolicies.hpp" +#include "src/include/mallocMC/AlignmentPolicies.hpp" // configurate the CreationPolicy "Scatter" @@ -71,10 +71,10 @@ struct AlignmentConfig{ // Define a new allocator and call it ScatterAllocator // which resembles the behaviour of ScatterAlloc -typedef PolicyMalloc::PolicyAllocator< - PolicyMalloc::CreationPolicies::Scatter, - PolicyMalloc::DistributionPolicies::XMallocSIMD, - PolicyMalloc::OOMPolicies::ReturnNull, - PolicyMalloc::ReservePoolPolicies::SimpleCudaMalloc, - PolicyMalloc::AlignmentPolicies::Shrink +typedef mallocMC::Allocator< + mallocMC::CreationPolicies::Scatter, + mallocMC::DistributionPolicies::XMallocSIMD, + mallocMC::OOMPolicies::ReturnNull, + mallocMC::ReservePoolPolicies::SimpleCudaMalloc, + mallocMC::AlignmentPolicies::Shrink > ScatterAllocator;