Commit 7f33dd8f authored by Martin Bauer's avatar Martin Bauer
Browse files

walberla boundary class now also works for GPU boundaries

parent c70e6432
......@@ -23,11 +23,14 @@
#include "core/DataTypes.h"
#include "core/Macros.h"
#include "{{class_name}}.h"
{% if target == 'gpu' -%}
#include "cuda/ErrorChecking.h"
{%- endif %}
{% if target is equalto 'cpu' -%}
{% if target == 'cpu' -%}
#define FUNC_PREFIX
{%- elif target is equalto 'gpu' -%}
{%- elif target == 'gpu' -%}
#define FUNC_PREFIX __global__
{%- endif %}
......@@ -39,6 +42,8 @@ namespace {{namespace}} {
#ifdef __GNUC__
#pragma GCC diagnostic push
#pragma GCC diagnostic ignored "-Wstrict-aliasing"
#pragma GCC diagnostic ignored "-Wunused-variable"
#pragma GCC diagnostic ignored "-Wconversion"
#endif
{{kernel|generate_definition}}
......@@ -50,15 +55,15 @@ namespace {{namespace}} {
void {{class_name}}::operator() ( IBlock * block )
{
auto * indexVector = block->getData<IndexVector>(indexVectorID);
auto * indexVectors = block->getData<IndexVectors>(indexVectorID);
auto & indexVector = indexVectors->{{target}};
int64_t indexVectorSize = int64_c( indexVectors->cpu.size() );
uint8_t * fd_indexVector = reinterpret_cast<uint8_t*>(&indexVector[0]);
int64_t indexVectorSize = int64_c( indexVector->size() );
uint8_t * fd_indexVector = reinterpret_cast<uint8_t*>(&(*indexVector)[0]);
{{kernel|generate_block_data_to_field_extraction(['indexVector', 'indexVectorSize'])|indent(4)}}
{{kernel|generate_call|indent(4)}}
{{kernel|generate_call(spatial_shape_symbols=['indexVectorSize'])|indent(4)}}
}
} // namespace {{namespace}}
} // namespace walberla
......@@ -50,14 +50,32 @@ class {{class_name}}
public:
{{StructDeclaration|indent(4)}}
typedef std::vector<{{StructName}}> IndexVector;
struct IndexVectors
{
bool operator==(IndexVectors & other) { return cpu == other.cpu; }
std::vector<{{StructName}}> cpu;
{% if target == 'gpu' -%}
{{StructName}} * gpu;
~IndexVectors() { cudaFree(gpu); }
IndexVectors() : gpu(nullptr) {}
void syncGPU()
{
cudaFree( gpu );
cudaMalloc( &gpu, sizeof({{StructName}}) * cpu.size() );
cudaMemcpy( gpu, &cpu[0], sizeof({{StructName}}) * cpu.size(), cudaMemcpyHostToDevice );
}
{% endif %}
};
{{class_name}}( const shared_ptr<StructuredBlockForest> & blocks,
{{kernel|generate_constructor_parameters(['indexVector', 'indexVectorSize'])}} )
: {{ kernel|generate_constructor_initializer_list(['indexVector', 'indexVectorSize']) }}
{
auto createIdxVector = []( IBlock * const , StructuredBlockStorage * const ) { return new IndexVector(); };
indexVectorID = blocks->addStructuredBlockData< IndexVector >( createIdxVector, "IndexField_{{class_name}}");
auto createIdxVector = []( IBlock * const , StructuredBlockStorage * const ) { return new IndexVectors(); };
indexVectorID = blocks->addStructuredBlockData< IndexVectors >( createIdxVector, "IndexField_{{class_name}}");
};
void operator() ( IBlock * block );
......@@ -75,13 +93,17 @@ public:
void fillFromFlagField( IBlock * block, ConstBlockDataID flagFieldID,
FlagUID boundaryFlagUID, FlagUID domainFlagUID )
{
auto * indexVector = block->getData< IndexVector > ( indexVectorID );
auto * indexVectors = block->getData< IndexVectors > ( indexVectorID );
auto & indexVector = indexVectors->cpu;
auto * flagField = block->getData< FlagField_T > ( flagFieldID );
auto boundaryFlag = flagField->getFlag(boundaryFlagUID);
auto domainFlag = flagField->getFlag(domainFlagUID);
indexVector->clear();
indexVector.clear();
{% if target == 'gpu' %}
cudaFree( indexVectors->gpu );
{% endif %}
for( auto it = flagField->begin(); it != flagField->end(); ++it )
{
......@@ -91,17 +113,22 @@ public:
{%- for dirIdx, offset in stencil_info %}
{% if dim == 2 -%}
if ( isFlagSet( it.neighbor({{offset}}, 0), boundaryFlag ) )
indexVector->push_back({{StructName}}(it.x(), it.y(), {{dirIdx}} ) );
indexVector.push_back({{StructName}}(it.x(), it.y(), {{dirIdx}} ) );
{%- elif dim == 3 -%}
if ( isFlagSet( it.neighbor({{offset}}), boundaryFlag ) )
indexVector->push_back({{StructName}}(it.x(), it.y(), it.z(), {{dirIdx}} ) );
indexVector.push_back({{StructName}}(it.x(), it.y(), it.z(), {{dirIdx}} ) );
{%- endif -%}
{% endfor %}
}
{% if target == 'gpu' %}
indexVectors->syncGPU();
{% endif %}
}
private:
BlockDataID indexVectorID;
{{kernel|generate_members(['indexVector', 'indexVectorSize'])|indent(4)}}
};
......
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