Commit d7963b9a authored by Tommaso Bianucci's avatar Tommaso Bianucci

Add binary search on keys and more structure in SparseGridGpu.

parent cecf8c37
This diff is collapsed.
......@@ -24,13 +24,22 @@ template<typename ScalarT>
struct DataBlock
{
ScalarT block[DATA_BLOCK_SIZE];
DataBlock(ScalarT backgroundValue = 0)
{
for (int i = 0; i < DATA_BLOCK_SIZE; ++i)
{
block[i] = backgroundValue;
}
}
virtual ScalarT &operator[](unsigned int index)
{
return block[index];
}
};
template <typename ScalarT>
template<typename ScalarT>
struct BackgroundDataBlock : DataBlock<ScalarT>
{
ScalarT backgroundValue = 0;
......@@ -53,7 +62,9 @@ private:
const unsigned int blocks = 1; // Template index for values
mgpu::standard_context_t context;
const unsigned int backgroundKey = 0;
ScalarT backgroundValue = 0;
openfpm::vector_gpu<aggregate<unsigned int, DataBlock<ScalarT>>> blockMap;
bool isSorted = true;
public:
SparseGridGpu(ScalarT backgroundValue, size_t initialCapacity = 1);
......@@ -65,14 +76,14 @@ public:
* @param linId The linear id of the element. Client must convert coordinates to linear id.
* @return The copy of the value at the given position.
*/
__device__ ScalarT value(unsigned int linId);
__host__ __device__ ScalarT value(unsigned int linId);
/**
* Get the value at a certain position of the sparse grid. This is for read-write use.
* @param linId The linear id of the element. Client must convert coordinates to linear id.
* @return The reference to the value at the given position.
*/
__device__ ScalarT &get(unsigned int linId);
__host__ __device__ ScalarT &get(unsigned int linId);
/**
* Synchronize data from host to device.
......@@ -88,6 +99,13 @@ private:
void resize(size_t newCapacity);
void setBackgroundValue(ScalarT backgroundValue);
/**
* Perform binary search on the keys of the blocks.
* @param key The key to search.
* @return The position of the key if found, 0 else.
*/
__host__ __device__ unsigned int binarySearchOnKeys(unsigned int key);
};
template<typename DataType>
......@@ -112,7 +130,9 @@ void SparseGridGpu<DataType>::resize(size_t newCapacity)
template<typename ScalarT>
void SparseGridGpu<ScalarT>::setBackgroundValue(ScalarT backgroundValue)
{
backgroundValue = backgroundValue;
blockMap.template get<blocks>(backgroundKey) = new BackgroundDataBlock<ScalarT>(backgroundValue);
//todo: here find a way to avoid this memory leak in case the background value needs to be reset for some reason
}
template<typename ScalarT>
......@@ -128,20 +148,66 @@ void SparseGridGpu<ScalarT>::deviceToHost()
}
template<typename ScalarT>
__device__ ScalarT SparseGridGpu<ScalarT>::value(unsigned int linId)
__host__ __device__ ScalarT SparseGridGpu<ScalarT>::value(unsigned int linId)
{
unsigned int bId = linId / DATA_BLOCK_SIZE;
unsigned int offset = linId % DATA_BLOCK_SIZE;
// Now perform binary search
unsigned int probePos = blockMap.size();
unsigned int probe = blockMap.template value<keys>(probePos);
while (bId != probe)
// Now perform binary search of the block
unsigned int resPos = binarySearchOnKeys(bId);
// Now we can retrieve the block and return the requested scalar value from it
// If the binary search failed, the returned value is background.
return blockMap.template get<blocks>(resPos)[offset];
}
template<typename ScalarT>
__host__ __device__ ScalarT &SparseGridGpu<ScalarT>::get(unsigned int linId)
{
unsigned int bId = linId / DATA_BLOCK_SIZE;
unsigned int offset = linId % DATA_BLOCK_SIZE;
// Now perform binary search of the block
unsigned int resPos = binarySearchOnKeys(bId);
// Now we can retrieve the block and return the requested scalar value from it
// If the binary search failed, however, we need to alloc a new block and return a reference to the new value.
if (resPos == 0)
{
//todo: here allocate a new block, set isSorted to false and return the reference to this value
DataBlock<ScalarT> &newBlock = new DataBlock<ScalarT>()
isSorted = false;
}
else
{
return blockMap.template get<blocks>(resPos)[offset];
}
}
template<typename ScalarT>
__host__ __device__ unsigned int SparseGridGpu<ScalarT>::binarySearchOnKeys(unsigned int key)
{
assert(isSorted); // We can binary search only if sorted
++key; // This is just an offset for the default block in position (and key) 0.
unsigned int a = 0, b = blockMap.size() - 1;
unsigned int mid = (a + b) / 2;
unsigned int probe = blockMap.template value<keys>(mid);
unsigned int resPos = 0;
while (a != b)
{
if (bId > probe)
if (key == probe)
{
resPos = mid;
break;
}
else if (key > probe)
{
a = mid;
}
else
{
probePos =
b = mid;
}
mid = (a + b) / 2;
probe = blockMap.template value<keys>(mid);
}
return resPos;
}
//template<typename DataType>
......
Markdown is supported
0% or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment