Skip to content

Fix small bugs #184

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
chengcli merged 4 commits into from
Jul 15, 2026
Merged

Fix small bugs #184

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
fef0a34
wip
Jul 15, 2026
d5d75d2
wip
Jul 15, 2026
789199b
wip
chengcli Jul 15, 2026
a105a58
wip
chengcli Jul 15, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
4 changes: 3 additions & 1 deletion src/output/load_diag_output_data.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -265,12 +265,14 @@ void OutputType::loadDiagOutputData(MeshBlockImpl* pmb, Variables const& vars) {
auto area = pcoord->face_area1();
auto ny = peos->nvar() - 5;
int il = pcoord->il();
int iu = pcoord->iu();

if (ny > 0) {
pod = new OutputData;
pod->type = "VECTORS";
pod->name = get_hydro_names(pmb, "path_");
auto u_sum = (u * vol).narrow(0, ICY, ny).sum(-1) / area.select(-1, il);
auto u_sum = (u * vol).narrow(0, ICY, ny).slice(-1, il, iu + 1).sum(-1) /
area.select(-1, il);

pod->data.CopyFromTensor(u_sum);
AppendOutputDataNode(pod);
Expand Down
2 changes: 1 addition & 1 deletion src/recon/cp3.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -11,7 +11,7 @@ namespace snap {

void Center3InterpImpl::reset() {
cm = register_buffer(
"cm", torch::tensor({-1. / 3., 5. / 6., -1. / 6.}, torch::kFloat64));
"cm", torch::tensor({1. / 3., 5. / 6., -1. / 6.}, torch::kFloat64));
cp = register_buffer("cp", cm.flip({0}));
}
Comment on lines 12 to 16

Expand Down
6 changes: 3 additions & 3 deletions src/recon/interp_impl.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -55,8 +55,8 @@ __device__ T interp_shared_weno3_coeff_impl(T const *line, T const *coeff,
T beta0 = SQR(_vvdot<3>(phi, c3));
T beta1 = SQR(_vvdot<3>(phi, c4));

T alpha0 = (1.0 / 3.0) / SQR(beta0 + 1e-6);
T alpha1 = (2.0 / 3.0) / SQR(beta1 + 1e-6);
T alpha0 = (2.0 / 3.0) / SQR(beta0 + 1e-6);
T alpha1 = (1.0 / 3.0) / SQR(beta1 + 1e-6);

return (alpha0 * p0 + alpha1 * p1) / (alpha0 + alpha1) * vscale;
}
Expand Down Expand Up @@ -109,7 +109,7 @@ __device__ T interp_shared_fused_impl(T const *line, int v, int start,
int axis_size, FusedReconScheme scheme,
bool right, bool scale) {
if (scheme == FusedReconScheme::CP3) {
constexpr T cm[3] = {-1. / 3., 5. / 6., -1. / 6.};
constexpr T cm[3] = {1. / 3., 5. / 6., -1. / 6.};
T c[3];
for (int k = 0; k < 3; ++k)
c[k] = right ? cm[2 - k] : cm[k];
Expand Down
44 changes: 19 additions & 25 deletions src/recon/interp_impl.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -10,7 +10,7 @@
namespace snap {

template <int N, typename T>
DISPATCH_MACRO inline T _vvdot(T *v1, T *v2) {
DISPATCH_MACRO inline T _vvdot(T* v1, T* v2) {
T out = 0.;
for (int i = 0; i < N; ++i) {
out += v1[i] * v2[i];
Expand All @@ -20,7 +20,7 @@ DISPATCH_MACRO inline T _vvdot(T *v1, T *v2) {

// polynomial
template <typename T, int N>
DISPATCH_MACRO void interp_poly_impl(T *out, T *inp, T *coeff, int stride1,
DISPATCH_MACRO void interp_poly_impl(T* out, T* inp, T* coeff, int stride1,
int stride2, int stride_out, int nvar) {
for (int j = 0; j < nvar; ++j) {
OUT(j) = 0.;
Expand All @@ -32,13 +32,13 @@ DISPATCH_MACRO void interp_poly_impl(T *out, T *inp, T *coeff, int stride1,

// WENO 3 interpolation
template <typename T>
DISPATCH_MACRO void interp_weno3_impl(T *out, T *inp, T *coeff, int stride1,
DISPATCH_MACRO void interp_weno3_impl(T* out, T* inp, T* coeff, int stride1,
int stride2, int stride_out, int nvar,
double scale) {
T *c1 = coeff;
T *c2 = c1 + 3;
T *c3 = c2 + 3;
T *c4 = c3 + 3;
T* c1 = coeff;
T* c2 = c1 + 3;
T* c3 = c2 + 3;
T* c4 = c3 + 3;

T phi[3];

Expand All @@ -56,39 +56,33 @@ DISPATCH_MACRO void interp_weno3_impl(T *out, T *inp, T *coeff, int stride1,
continue;
}

if (vscale != 0.0) {
phi[0] /= vscale;
phi[1] /= vscale;
phi[2] /= vscale;
}

T p0 = _vvdot<3>(phi, c1);
T p1 = _vvdot<3>(phi, c2);

T beta0 = SQR(_vvdot<3>(phi, c3));
T beta1 = SQR(_vvdot<3>(phi, c4));

T alpha0 = (1.0 / 3.0) / SQR(beta0 + 1e-6);
T alpha1 = (2.0 / 3.0) / SQR(beta1 + 1e-6);
T alpha0 = (2.0 / 3.0) / SQR(beta0 + 1e-6);
T alpha1 = (1.0 / 3.0) / SQR(beta1 + 1e-6);

OUT(j) = (alpha0 * p0 + alpha1 * p1) / (alpha0 + alpha1) * vscale;
}
};

// WENO 5 interpolation
template <typename T>
DISPATCH_MACRO void interp_weno5_impl(T *out, T *inp, T *coeff, int stride1,
DISPATCH_MACRO void interp_weno5_impl(T* out, T* inp, T* coeff, int stride1,
int stride2, int stride_out, int nvar,
double scale) {
T *c1 = coeff;
T *c2 = c1 + 5;
T *c3 = c2 + 5;
T *c4 = c3 + 5;
T *c5 = c4 + 5;
T *c6 = c5 + 5;
T *c7 = c6 + 5;
T *c8 = c7 + 5;
T *c9 = c8 + 5;
T* c1 = coeff;
T* c2 = c1 + 5;
T* c3 = c2 + 5;
T* c4 = c3 + 5;
T* c5 = c4 + 5;
T* c6 = c5 + 5;
T* c7 = c6 + 5;
T* c8 = c7 + 5;
T* c9 = c8 + 5;

T phi[5];

Expand Down
4 changes: 2 additions & 2 deletions src/recon/interp_simple.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -99,8 +99,8 @@ inline DISPATCH_MACRO T interp_weno3(T const& phim1, T const& phi,
T beta0 = (phim1 - phi) * (phim1 - phi);
T beta1 = (phi - phip1) * (phi - phip1);

T alpha0 = (1.0 / 3.0) / sqr(beta0 + 1e-6);
T alpha1 = (2.0 / 3.0) / sqr(beta1 + 1e-6);
T alpha0 = (2.0 / 3.0) / sqr(beta0 + 1e-6);
T alpha1 = (1.0 / 3.0) / sqr(beta1 + 1e-6);

return (alpha0 * p0 + alpha1 * p1) / (alpha0 + alpha1);
};
Expand Down
4 changes: 2 additions & 2 deletions src/recon/recon_dispatch.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -87,8 +87,8 @@ void call_weno3_mps(at::TensorIterator& iter, torch::Tensor coeff, int dim,
wu /= wscale.unsqueeze(-1);
}

auto alpha1 = 1. / 3. / (wu.matmul(c3).square() + 1e-6).square();
auto alpha2 = 2. / 3. / (wu.matmul(c4).square() + 1e-6).square();
auto alpha1 = 2. / 3. / (wu.matmul(c3).square() + 1e-6).square();
auto alpha2 = 1. / 3. / (wu.matmul(c4).square() + 1e-6).square();

torch::add_out(result, alpha1 * wu.matmul(c1), alpha2 * wu.matmul(c2));
result /= alpha1 + alpha2;
Expand Down
8 changes: 7 additions & 1 deletion tests/run_shallow_splash.cmake
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -17,7 +17,13 @@ if(NOT _status EQUAL 0)
message(FATAL_ERROR "Failed to download reference file with exit code ${_status}")
endif()

execute_process(COMMAND ln -sf ../bin/shallow_splash.yaml shallow_splash.yaml)
file(READ "../bin/shallow_splash.yaml" config)
file(READ "../configure.h" configure_h)
if(configure_h MATCHES "NO_PNETCDFOUTPUT")
string(REPLACE "type: pnetcdf" "type: netcdf" config "${config}")
endif()
file(REMOVE "shallow_splash.yaml")
file(WRITE "shallow_splash.yaml" "${config}")

execute_process(
COMMAND torchrun --no-python --nproc-per-node=6 ../bin/shallow_splash.${buildl}
Expand Down
8 changes: 7 additions & 1 deletion tests/run_straka.cmake
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -17,7 +17,13 @@ if(NOT _status EQUAL 0)
message(FATAL_ERROR "Failed to download reference file with exit code ${_status}")
endif()

execute_process(COMMAND ln -sf ../bin/straka.yaml straka.yaml)
file(READ "../bin/straka.yaml" config)
file(READ "../configure.h" configure_h)
if(configure_h MATCHES "NO_PNETCDFOUTPUT")
string(REPLACE "type: pnetcdf" "type: netcdf" config "${config}")
endif()
file(REMOVE "straka.yaml")
file(WRITE "straka.yaml" "${config}")

execute_process(
COMMAND torchrun --no-python --nproc-per-node=2 ../bin/straka.${buildl}
Expand Down
50 changes: 29 additions & 21 deletions tests/test_fused_cs_sync_smoke.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -47,20 +47,6 @@ torch::Device select_cuda_device() {
return torch::Device(torch::kCUDA, local_rank);
}

void require_cuda_6rank_or_skip() {
if (!torch::cuda::is_available()) {
GTEST_SKIP() << "CUDA runtime is unavailable";
}
int world_size = env_int("WORLD_SIZE", 1);
if (world_size != 6) {
GTEST_SKIP() << "test_fused_cs_sync_smoke requires torchrun with 6 ranks";
}
int device_count = c10::cuda::device_count();
if (device_count < world_size) {
GTEST_SKIP() << "test_fused_cs_sync_smoke requires 6 CUDA devices";
}
}

LayoutOptions make_layout_options() {
auto opts = LayoutOptionsImpl::create();
opts->type("cubed-sphere");
Expand Down Expand Up @@ -134,6 +120,21 @@ bool fused_exchange_uses_right_interp_for_side(int side) {
return !fused_exchange_uses_right_state_start_for_side(side);
}

bool cuda_6rank_available_or_skip() {
if (!torch::cuda::is_available()) {
return false;
}
int world_size = env_int("WORLD_SIZE", 1);
if (world_size != 6) {
return false;
}
int device_count = c10::cuda::device_count();
if (device_count < world_size) {
return false;
}
return true;
}

void initialize_symmetric_memory_group(LayoutImpl const &layout,
std::string const &group_name) {
static std::set<std::string> initialized;
Expand Down Expand Up @@ -318,7 +319,8 @@ __global__ void verify_hydro_remote_constant_kernel(void **buffer_ptrs,
} // namespace

TEST(FusedCubedSphereSymmetricMemory, RendezvousCompletes) {
require_cuda_6rank_or_skip();
if (!cuda_6rank_available_or_skip())
GTEST_SKIP();
auto ctx = make_smoke_context();
std::string group_name = "snapy:test:fused-cs-rendezvous";
initialize_symmetric_memory_group(*ctx.layout, group_name);
Expand All @@ -331,7 +333,8 @@ TEST(FusedCubedSphereSymmetricMemory, RendezvousCompletes) {
}

TEST(FusedCubedSphereSymmetricMemory, PreviousKernelSyncCompletes) {
require_cuda_6rank_or_skip();
if (!cuda_6rank_available_or_skip())
GTEST_SKIP();
auto ctx = make_smoke_context();
std::string group_name = "snapy:test:fused-cs-sync";
initialize_symmetric_memory_group(*ctx.layout, group_name);
Expand All @@ -357,7 +360,8 @@ TEST(FusedCubedSphereSymmetricMemory, PreviousKernelSyncCompletes) {
}

TEST(FusedCubedSphereSymmetricMemory, OrientationMetadataMatchesStaged) {
require_cuda_6rank_or_skip();
if (!cuda_6rank_available_or_skip())
GTEST_SKIP();
auto ctx = make_smoke_context();

auto meta = ctx.edge_meta.cpu();
Expand All @@ -378,7 +382,8 @@ TEST(FusedCubedSphereSymmetricMemory, OrientationMetadataMatchesStaged) {
}

TEST(FusedCubedSphereSymmetricMemory, BoundaryStateConventionMatchesStaged) {
require_cuda_6rank_or_skip();
if (!cuda_6rank_available_or_skip())
GTEST_SKIP();
// Staged hydro sends the local right state to lower-side neighbors and the
// local left state to upper-side neighbors.
EXPECT_TRUE(fused_exchange_uses_right_state_start_for_side(SIDE_L));
Expand All @@ -392,7 +397,8 @@ TEST(FusedCubedSphereSymmetricMemory, BoundaryStateConventionMatchesStaged) {
}

TEST(FusedCubedSphereSymmetricMemory, RemoteEdgePayloadMatches) {
require_cuda_6rank_or_skip();
if (!cuda_6rank_available_or_skip())
GTEST_SKIP();
auto ctx = make_smoke_context();
std::string group_name = "snapy:test:fused-cs-remote-read";
initialize_symmetric_memory_group(*ctx.layout, group_name);
Expand Down Expand Up @@ -427,7 +433,8 @@ TEST(FusedCubedSphereSymmetricMemory, RemoteEdgePayloadMatches) {
}

TEST(FusedCubedSphereSymmetricMemory, RemoteStateSelectionMatchesStaged) {
require_cuda_6rank_or_skip();
if (!cuda_6rank_available_or_skip())
GTEST_SKIP();
auto ctx = make_smoke_context();
std::string group_name = "snapy:test:fused-cs-remote-state";
initialize_symmetric_memory_group(*ctx.layout, group_name);
Expand Down Expand Up @@ -463,7 +470,8 @@ TEST(FusedCubedSphereSymmetricMemory, RemoteStateSelectionMatchesStaged) {
}

TEST(FusedCubedSphereSymmetricMemory, ConstantStateExchangeHasZeroMassFlux) {
require_cuda_6rank_or_skip();
if (!cuda_6rank_available_or_skip())
GTEST_SKIP();
auto ctx = make_smoke_context();
std::string group_name = "snapy:test:fused-cs-constant-state";
initialize_symmetric_memory_group(*ctx.layout, group_name);
Expand Down
Loading
Loading