Skip to content

Commit d977704

Browse files
tobre1tobre1
authored
Double precision GPU result buffer (#186)
* Implement double precision result buffer on GPU * Bump deps
1 parent c9561a7 commit d977704

15 files changed

Lines changed: 73 additions & 52 deletions

‎CMakeLists.txt‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -101,7 +101,7 @@ include("cmake/cpm.cmake")
101101

102102
CPMAddPackage(
103103
NAME ViennaCore
104-
VERSION 1.7.3
104+
VERSION 1.7.4
105105
GIT_REPOSITORY "https://github.com/ViennaTools/ViennaCore"
106106
EXCLUDE_FROM_ALL ${VIENNAPS_BUILD_PYTHON}
107107
OPTIONS "VIENNACORE_USE_GPU ${VIENNAPS_USE_GPU}")
@@ -113,7 +113,7 @@ CPMAddPackage(
113113

114114
CPMFindPackage(
115115
NAME ViennaRay
116-
VERSION 3.8.5
116+
VERSION 3.9.0
117117
GIT_REPOSITORY "https://github.com/ViennaTools/ViennaRay"
118118
EXCLUDE_FROM_ALL ${VIENNAPS_BUILD_PYTHON}
119119
OPTIONS "VIENNARAY_USE_GPU ${VIENNAPS_USE_GPU}")

‎cmake/generate_ptx.cmake‎

Lines changed: 10 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -12,6 +12,9 @@ function(generate_pipeline target_name generated_files_output)
1212
cuda_include_directories(${ViennaPS_SOURCE_DIR}/include/viennaps)
1313
cuda_include_directories(${ViennaCore_SOURCE_DIR}/include/viennacore)
1414
add_compile_definitions(VIENNACORE_COMPILE_GPU)
15+
if(VIENNARAY_GPU_DOUBLE_PRECISION)
16+
add_compile_definitions(VIENNARAY_GPU_DOUBLE_PRECISION)
17+
endif()
1518

1619
# Generate OptiX IR files if enabled
1720
if(VIENNAPS_GENERATE_OPTIXIR)
@@ -54,6 +57,9 @@ function(generate_kernel target_name generated_files_output)
5457
cuda_include_directories(${ViennaCore_SOURCE_DIR}/include/viennacore)
5558
cuda_include_directories(${VIENNARAY_GPU_INCLUDE} ${VIENNAPS_GPU_INCLUDE})
5659
add_compile_definitions(VIENNACORE_COMPILE_GPU)
60+
if(VIENNARAY_GPU_DOUBLE_PRECISION)
61+
add_compile_definitions(VIENNARAY_GPU_DOUBLE_PRECISION)
62+
endif()
5763

5864
cuda_wrap_srcs(
5965
${target_name}
@@ -100,6 +106,9 @@ function(add_GPU_executable target_name_base target_name_var)
100106
cuda_include_directories(${ViennaPS_SOURCE_DIR}/include/viennaps)
101107
cuda_include_directories(${ViennaCore_SOURCE_DIR}/include/viennacore)
102108
add_compile_definitions(VIENNACORE_COMPILE_GPU)
109+
if(VIENNARAY_GPU_DOUBLE_PRECISION)
110+
add_compile_definitions(VIENNARAY_GPU_DOUBLE_PRECISION)
111+
endif()
103112

104113
# Create CUDA kernels
105114
cuda_wrap_srcs(
@@ -145,7 +154,7 @@ function(add_GPU_executable target_name_base target_name_var)
145154
# Here is where we create the rule to make the executable. We define a target name and
146155
# list all the source files used to create the target. In addition we also pass along
147156
# the cmake_options parsed out of the arguments.
148-
message(STATUS "Adding target: ${target_name}")
157+
message(STATUS "Adding GPU target: ${target_name}")
149158
add_executable(${target_name} ${source_files} ${generated_files} ${cmake_options})
150159
target_include_directories(${target_name} PRIVATE ${VIENNARAY_GPU_INCLUDE}
151160
${VIENNAPS_GPU_INCLUDE})

‎gpu/benchmark/CPU_Benchmark.cpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -186,7 +186,7 @@ int main() {
186186
tracer.normalizeFlux(fluxResult);
187187
fluxResultVec.push_back(std::move(fluxResult));
188188
}
189-
ElementToPointData<NumericType, float>(
189+
ElementToPointData<NumericType, float, viennaray::gpu::ResultType>(
190190
fluxResultVec, pointData, particles, elementKdTree, diskMesh,
191191
surfMesh, domain->getGridDelta() * 2.0f)
192192
.apply();

‎gpu/benchmark/GPU_Benchmark.cpp‎

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -111,7 +111,7 @@ int main() {
111111
auto pointData = viennals::PointData<NumericType>::New();
112112
tracer.normalizeResults();
113113
// tracer.downloadResults();
114-
gpu::ElementToPointData<NumericType, float>(
114+
gpu::ElementToPointData<NumericType, float, viennaray::gpu::ResultType>(
115115
tracer.getResults(), pointData, tracer.getParticles(),
116116
elementKdTree, diskMesh, surfMesh, domain->getGridDelta() * 2.0f)
117117
.apply();
@@ -211,8 +211,9 @@ int main() {
211211
tracer.downloadResults();
212212
int smoothingNeighbors = 1;
213213
auto flux = tracer.getFlux(0, 0, smoothingNeighbors);
214+
std::vector<NumericType> fluxNumeric(flux.begin(), flux.end());
214215
auto velocities =
215-
SmartPointer<std::vector<NumericType>>::New(std::move(flux));
216+
SmartPointer<std::vector<NumericType>>::New(std::move(fluxNumeric));
216217
velocityField->prepare(domain, velocities, 0.);
217218
timer.finish();
218219
file << timer.currentDuration << ";";

‎gpu/benchmark/GPU_Benchmark_single.cpp‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -95,9 +95,9 @@ int main() {
9595
auto flux = tracer.getFlux(0, 0);
9696
surfMesh->getCellData().insertNextScalarData(flux, "flux");
9797
auto pointData = viennals::PointData<NumericType>::New();
98-
gpu::ElementToPointData<NumericType>(tracer.getResults(), pointData,
99-
tracer.getParticles(), elementKdTree,
100-
diskMesh, surfMesh, GRID_DELTA)
98+
gpu::ElementToPointData<NumericType, float, viennaray::gpu::ResultType>(
99+
tracer.getResults(), pointData, tracer.getParticles(), elementKdTree,
100+
diskMesh, surfMesh, GRID_DELTA)
101101
.apply();
102102
timer.finish();
103103
file << timer.currentDuration;

‎gpu/benchmark/compare_flux.cpp‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -83,9 +83,9 @@ int main() {
8383
tracer.apply();
8484

8585
auto pointData = viennals::PointData<NumericType>::New();
86-
gpu::ElementToPointData<NumericType>(tracer.getResults(), pointData,
87-
tracer.getParticles(), elementKdTree,
88-
diskMesh, surfMesh, GRID_DELTA)
86+
gpu::ElementToPointData<NumericType, float, viennaray::gpu::ResultType>(
87+
tracer.getResults(), pointData, tracer.getParticles(), elementKdTree,
88+
diskMesh, surfMesh, GRID_DELTA)
8989
.apply();
9090

9191
downloadResultsToPointData(tracer, surfMesh->getCellData());

‎gpu/include/models/FaradayCage.cuh‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -41,13 +41,13 @@ faradayIonCollision(const void *sbtData, viennaray::gpu::PerRayData *prd) {
4141

4242
// flux array
4343
atomicAdd(&launchParams.resultBuffer[prd->primIDs[i]],
44-
prd->rayWeight * yield);
44+
(viennaray::gpu::ResultType)prd->rayWeight * yield);
4545

4646
if (params->redepositionRate > 0.f) {
4747
// redeposition array
4848
atomicAdd(&launchParams
4949
.resultBuffer[launchParams.numElements + prd->primIDs[i]],
50-
prd->load);
50+
(viennaray::gpu::ResultType)prd->load);
5151
}
5252
}
5353
}

‎gpu/include/models/IonBeamEtching.cuh‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -41,13 +41,13 @@ __forceinline__ __device__ void IBECollision(const void *sbtData,
4141
// flux array
4242
atomicAdd(&launchParams.resultBuffer[getIdxOffset(0, launchParams) +
4343
prd->primIDs[i]],
44-
prd->rayWeight * yield);
44+
(viennaray::gpu::ResultType)prd->rayWeight * yield);
4545

4646
if (redepositionEnabled) {
4747
// redeposition array
4848
atomicAdd(&launchParams.resultBuffer[getIdxOffset(1, launchParams) +
4949
prd->primIDs[i]],
50-
prd->load);
50+
(viennaray::gpu::ResultType)prd->load);
5151
}
5252
}
5353
}

‎gpu/include/models/MultiParticle.cuh‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -20,7 +20,7 @@ multiNeutralCollision(viennaray::gpu::PerRayData *prd) {
2020
atomicAdd(&launchParams
2121
.resultBuffer[viennaray::gpu::getIdxOffset(0, launchParams) +
2222
prd->primIDs[i]],
23-
prd->rayWeight);
23+
(viennaray::gpu::ResultType)prd->rayWeight);
2424
}
2525
}
2626

@@ -59,7 +59,7 @@ multiIonCollision(const void *sbtData, viennaray::gpu::PerRayData *prd) {
5959
atomicAdd(&launchParams
6060
.resultBuffer[viennaray::gpu::getIdxOffset(0, launchParams) +
6161
prd->primIDs[i]],
62-
flux);
62+
(viennaray::gpu::ResultType)flux);
6363
}
6464
}
6565

‎gpu/include/models/PlasmaEtching.cuh‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -22,7 +22,7 @@ plasmaNeutralCollision(viennaray::gpu::PerRayData *prd) {
2222
atomicAdd(&launchParams
2323
.resultBuffer[viennaray::gpu::getIdxOffset(0, launchParams) +
2424
prd->primIDs[i]],
25-
prd->rayWeight);
25+
(viennaray::gpu::ResultType)prd->rayWeight);
2626
}
2727
}
2828

0 commit comments

Comments
 (0)