Skip to content

Feat/kurganov shared memory #117

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

Open
wants to merge 38 commits into
base: development
Choose a base branch
from
Open

Feat/kurganov shared memory #117

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
38 commits
Select commit Hold shift + click to select a range
3e17570
add structure for parallel process of rivers
CyprienBosserelle Sep 24, 2024
5918ee6
try to build new array to parallellise rivers
CyprienBosserelle Sep 25, 2024
3011e4d
Fix some compile issue
CyprienBosserelle Sep 25, 2024
65c82b6
Fix compile issue
CyprienBosserelle Sep 25, 2024
6c7489d
Fix code and fill Xriver array
CyprienBosserelle Sep 25, 2024
c5c3cbe
fix allocation blunder
CyprienBosserelle Oct 3, 2024
e058769
Add functions for Pinned memory
CyprienBosserelle Oct 10, 2024
9cb83cb
Fix Compile paged mem
CyprienBosserelle Oct 10, 2024
1b6c499
Make test for paged mem
CyprienBosserelle Oct 11, 2024
9fa3614
tweak test and remove out to screen
CyprienBosserelle Oct 13, 2024
00f02f3
modify test for Pin Meme for non-GPU
CyprienBosserelle Oct 15, 2024
9c59e3e
Fix CPU only MappedMem
CyprienBosserelle Oct 16, 2024
236b679
add allocations of river info in GPU XModel
CyprienBosserelle Oct 16, 2024
b8e8ed9
Add missing variable but also added template to various classes
CyprienBosserelle Oct 17, 2024
ef90c69
Fix compile issues
CyprienBosserelle Oct 17, 2024
31ab89f
fix map mem alloc
CyprienBosserelle Oct 17, 2024
b0576b1
Add momentum adjustment when using rain
CyprienBosserelle Oct 18, 2024
a957e97
Add limiter for flux adjustment for dry cells
CyprienBosserelle Oct 22, 2024
1c7018f
playing up with velocity sanity
CyprienBosserelle Oct 29, 2024
56bfaac
Add explanation to new algo
CyprienBosserelle Oct 29, 2024
306cc11
Revert experiment changes on roughness
CyprienBosserelle Jan 8, 2025
bda4751
Fix Dynamic forcing
CyprienBosserelle Jan 16, 2025
a835af5
force sane velocity
CyprienBosserelle Jan 30, 2025
9ba730a
Fix zsoffset
CyprienBosserelle Feb 4, 2025
ad71ffa
Clean <<< ...>>>
CyprienBosserelle Feb 11, 2025
db30cd0
Adding record of timestart
CyprienBosserelle Mar 6, 2025
8aa4c77
Update Makefile
CyprienBosserelle Mar 10, 2025
50c7238
ad bnd filter and relax time to input param
CyprienBosserelle Mar 11, 2025
385fcdc
apply tapering on non-uniform bnd
CyprienBosserelle Apr 28, 2025
e499aad
Update InitialConditions.cu
CyprienBosserelle Apr 28, 2025
2942cc3
add "all" and "file" as keyword for boundary forcing side
CyprienBosserelle Apr 29, 2025
83ae6ad
add sm50 to Makefile for auto test
May 22, 2025
a4c4094
Update Param.h
CyprienBosserelle May 22, 2025
8ae704d
add switch to enforce mass conservation in timestep selection
CyprienBosserelle May 23, 2025
cc07697
add mass conservation forcing as switch to param file
CyprienBosserelle May 23, 2025
55771dc
Update ReadInput.cu
CyprienBosserelle May 24, 2025
530c046
Update Param.h
CyprienBosserelle May 24, 2025
de7147a
I've made some improvements to the CUDA kernels in `src/Kurganov.cu` …
google-labs-jules[bot] May 25, 2025
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
13 changes: 8 additions & 5 deletions Makefile
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -198,7 +198,7 @@ ALL_CCFLAGS += $(addprefix -Xcompiler ,$(EXTRA_CCFLAGS))

SAMPLE_ENABLED := 1

ALL_LDFLAGS := -lnetcdf -I
ALL_LDFLAGS := -I
ALL_LDFLAGS += $(ALL_CCFLAGS)
ALL_LDFLAGS += $(addprefix -Xlinker ,$(LDFLAGS))
ALL_LDFLAGS += $(addprefix -Xlinker ,$(EXTRA_LDFLAGS))
Expand All @@ -207,10 +207,14 @@ ALL_LDFLAGS += $(addprefix -Xlinker ,$(EXTRA_LDFLAGS))
INCLUDES := -I/usr/includes
LIBRARIES :=

# Add NetCDF include library
INCLUDES += $(shell nc-config --cflags)
ALL_LDFLAGS += $(shell nc-config --libs)

################################################################################

# Gencode arguments
SMS ?= 35 50 52 60
SMS ?= 50 52 60 75
#SMS ?= 20 30 35
ifeq ($(SMS),)
$(info >>> WARNING - no SM architectures have been specified - waiving sample <<<)
Expand Down Expand Up @@ -318,7 +322,7 @@ ReadInput.o:./src/ReadInput.cu

Read_netcdf.o:./src/Read_netcdf.cu
$(EXEC) $(NVCC) $(INCLUDES) $(ALL_CCFLAGS) $(GENCODE_FLAGS) -o $@ -c $<

Reimann.o:./src/Reimann.cu
$(EXEC) $(NVCC) $(INCLUDES) $(ALL_CCFLAGS) $(GENCODE_FLAGS) -o $@ -c $<

Expand All @@ -338,8 +342,7 @@ Write_netcdf.o:./src/Write_netcdf.cu
$(EXEC) $(NVCC) $(INCLUDES) $(ALL_CCFLAGS) $(GENCODE_FLAGS) -o $@ -c $<

Write_txtlog.o:./src/Write_txtlog.cpp
$(EXEC) $(NVCC) $(INCLUDES) $(ALL_CCFLAGS) $(GENCODE_FLAGS) -o $@ -c $<

$(EXEC) $(NVCC) $(INCLUDES) $(ALL_CCFLAGS) $(GENCODE_FLAGS) -o $@ -c $<

Spherical.o:./src/Spherical.cu
$(EXEC) $(NVCC) $(INCLUDES) $(ALL_CCFLAGS) $(GENCODE_FLAGS) -o $@ -c $<
Expand Down
14 changes: 7 additions & 7 deletions src/Advection.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -251,11 +251,11 @@ template <class T> __global__ void AdvkernelGPU(Param XParam, BlockP<T> XBlock,
T ho, uo, vo;
T dhi = XAdv.dh[i];

T edt = dt;// dhi > T(0.0) ? dt : min(dt, hold / (T(-1.0) * dhi));

T edt = XParam.ForceMassConserve ? dt : dhi >= T(0.0) ? dt : min(dt, max(hold, XParam.eps) / abs(dhi));

//ho = max(hold + edt * dhi,T(0.0));
ho = hold + edt * dhi;


if (ho > eps) {
//
uo = (hold * XEv.u[i] + edt * XAdv.dhu[i]) / ho;
Expand Down Expand Up @@ -307,7 +307,7 @@ template <class T> __host__ void AdvkernelCPU(Param XParam, BlockP<T> XBlock, T

dhi = XAdv.dh[i];

T edt = dt;// dhi > T(0.0) ? dt : min(dt, hold / (T(-1.0) * dhi));
T edt = XParam.ForceMassConserve ? dt : dhi >= T(0.0) ? dt : min(dt, max(hold, XParam.eps) / abs(dhi));

ho = hold + edt * dhi;

Expand Down Expand Up @@ -475,7 +475,7 @@ template <class T> __host__ T CalctimestepGPU(Param XParam,Loop<T> XLoop, BlockP

//GPU Harris reduction #3. 8.3x reduction #0 Note #7 if a lot faster
// This was successfully tested with a range of grid size
//reducemax3 << <gridDimLine, blockDimLine, 64*sizeof(float) >> >(dtmax_g, arrmax_g, nx*ny)
//reducemax3 <<<gridDimLine, blockDimLine, 64*sizeof(float) >>>(dtmax_g, arrmax_g, nx*ny)

int maxThreads = 256;
int threads = (s < maxThreads * 2) ? nextPow2((s + 1) / 2) : maxThreads;
Expand All @@ -485,7 +485,7 @@ template <class T> __host__ T CalctimestepGPU(Param XParam,Loop<T> XLoop, BlockP
dim3 gridDimLine(blocks, 1, 1);


reducemin3 << <gridDimLine, blockDimLine, smemSize >> > (XTime.dtmax, XTime.arrmin, s);
reducemin3 <<<gridDimLine, blockDimLine, smemSize >>> (XTime.dtmax, XTime.arrmin, s);
CUDA_CHECK(cudaDeviceSynchronize());


Expand All @@ -503,7 +503,7 @@ template <class T> __host__ T CalctimestepGPU(Param XParam,Loop<T> XLoop, BlockP

CUDA_CHECK(cudaMemcpy(XTime.dtmax, XTime.arrmin, s * sizeof(T), cudaMemcpyDeviceToDevice));

reducemin3 << <gridDimLineS, blockDimLineS, smemSize >> > (XTime.dtmax, XTime.arrmin, s);
reducemin3 <<<gridDimLineS, blockDimLineS, smemSize >>> (XTime.dtmax, XTime.arrmin, s);
CUDA_CHECK(cudaDeviceSynchronize());

s = (s + (threads * 2 - 1)) / (threads * 2);
Expand Down
28 changes: 24 additions & 4 deletions src/Arrays.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -82,6 +82,23 @@ struct maskinfo

};

template <class T>
struct RiverInfo
{
int nbir;
int nburmax; // size of (max number of) unique block with rivers
int nribmax; // size of (max number of) rivers in one block
int* Xbidir; // array of block id for each river size(nburmax,nribmax)
int* Xridib; // array of river id in each block size(nburmax,nribmax)
T* xstart;
T* xend;
T* ystart;
T *yend;
T* qnow; // qnow is a pin mapped and so both pointers are needed here
T* qnow_g; // this simplify the code later

};


// outzone info used to actually write the nc files (one nc file by zone, the default zone is the full domain)
struct outzoneB
Expand Down Expand Up @@ -125,7 +142,7 @@ struct AdaptP




template <class T>
struct BndblockP
{
int nblkriver, nblkTs, nbndblkleft, nbndblkright, nbndblktop, nbndblkbot;
Expand All @@ -140,12 +157,15 @@ struct BndblockP
int* top;
int* bot;


RiverInfo<T> Riverinfo;


};


struct RiverBlk
{
std::vector<int> block;
};



Expand Down Expand Up @@ -208,7 +228,7 @@ struct Model

AdaptP adapt;

BndblockP bndblk;
BndblockP<T> bndblk;



Expand Down
27 changes: 16 additions & 11 deletions src/Boundary.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -95,6 +95,11 @@ template <class T> void FlowbndFlux(Param XParam, double totaltime, BlockP<T> XB
T taper=T(1.0);
if (bndseg.on)
{
if (XParam.bndtaper > 0.0)
{
taper = min((totaltime - XParam.inittime) / XParam.bndtaper, 1.0);
}

if (bndseg.uniform)
{
int SLstepinbnd = 1;
Expand All @@ -110,10 +115,7 @@ template <class T> void FlowbndFlux(Param XParam, double totaltime, BlockP<T> XB
zsbnd = interptime(bndseg.data[SLstepinbnd].wspeed, bndseg.data[SLstepinbnd - 1].wspeed, bndseg.data[SLstepinbnd].time - bndseg.data[SLstepinbnd - 1].time, totaltime - bndseg.data[SLstepinbnd - 1].time);


if (XParam.bndtaper > 0.0)
{
taper = min(totaltime / XParam.bndtaper, 1.0);
}

}
else
{
Expand All @@ -131,25 +133,25 @@ template <class T> void FlowbndFlux(Param XParam, double totaltime, BlockP<T> XB
//Left
//template <class T> __global__ void bndFluxGPUSide(Param XParam, bndsegmentside side, BlockP<T> XBlock, DynForcingP<float> Atmp, DynForcingP<float> Zsmap, bool uniform, float zsbnd, T * zs, T * h, T * un, T * ut, T * Fh, T * Fq, T * Ss)
//bndFluxGPUSide <<< gridDimBBND, blockDim, 0 >>> (XParam, bndseg.left, XBlock, Atmp, bndseg.WLmap, bndseg.uniform, bndseg.type, float(zsbnd), XEv.zs, XEv.h, un, ut, Fh, Fq, S);
bndFluxGPUSide << < gridDimBBNDLeft, blockDim, 0 >> > (XParam, bndseg.left, XBlock, Atmp, bndseg.WLmap, bndseg.uniform, bndseg.type, float(zsbnd), taper, XEv.zs, XEv.h, XEv.u, XEv.v, XFlux.Fhu, XFlux.Fqux, XFlux.Su);
bndFluxGPUSide <<< gridDimBBNDLeft, blockDim, 0 >>> (XParam, bndseg.left, XBlock, Atmp, bndseg.WLmap, bndseg.uniform, bndseg.type, float(zsbnd), taper, XEv.zs, XEv.h, XEv.u, XEv.v, XFlux.Fhu, XFlux.Fqux, XFlux.Su);
CUDA_CHECK(cudaDeviceSynchronize());
}
//if (bndseg.right.nblk > 0)
{
//Right
bndFluxGPUSide << < gridDimBBNDRight, blockDim, 0 >> > (XParam, bndseg.right, XBlock, Atmp, bndseg.WLmap, bndseg.uniform, bndseg.type, float(zsbnd), taper, XEv.zs, XEv.h, XEv.u, XEv.v, XFlux.Fhu, XFlux.Fqux, XFlux.Su);
bndFluxGPUSide <<< gridDimBBNDRight, blockDim, 0 >>> (XParam, bndseg.right, XBlock, Atmp, bndseg.WLmap, bndseg.uniform, bndseg.type, float(zsbnd), taper, XEv.zs, XEv.h, XEv.u, XEv.v, XFlux.Fhu, XFlux.Fqux, XFlux.Su);
CUDA_CHECK(cudaDeviceSynchronize());
}
//if (bndseg.top.nblk > 0)
{
//top
bndFluxGPUSide << < gridDimBBNDTop, blockDim, 0 >> > (XParam, bndseg.top, XBlock, Atmp, bndseg.WLmap, bndseg.uniform, bndseg.type, float(zsbnd), taper, XEv.zs, XEv.h, XEv.v, XEv.u, XFlux.Fhv, XFlux.Fqvy, XFlux.Sv);
bndFluxGPUSide <<< gridDimBBNDTop, blockDim, 0 >>> (XParam, bndseg.top, XBlock, Atmp, bndseg.WLmap, bndseg.uniform, bndseg.type, float(zsbnd), taper, XEv.zs, XEv.h, XEv.v, XEv.u, XFlux.Fhv, XFlux.Fqvy, XFlux.Sv);
CUDA_CHECK(cudaDeviceSynchronize());
}
//if (bndseg.bot.nblk > 0)
{
//bot
bndFluxGPUSide << < gridDimBBNDBot, blockDim, 0 >> > (XParam, bndseg.bot, XBlock, Atmp, bndseg.WLmap, bndseg.uniform, bndseg.type, float(zsbnd), taper, XEv.zs, XEv.h, XEv.v, XEv.u, XFlux.Fhv, XFlux.Fqvy, XFlux.Sv);
bndFluxGPUSide <<< gridDimBBNDBot, blockDim, 0 >>> (XParam, bndseg.bot, XBlock, Atmp, bndseg.WLmap, bndseg.uniform, bndseg.type, float(zsbnd), taper, XEv.zs, XEv.h, XEv.v, XEv.u, XFlux.Fhv, XFlux.Fqvy, XFlux.Sv);
CUDA_CHECK(cudaDeviceSynchronize());
}
}
Expand Down Expand Up @@ -219,7 +221,7 @@ template <class T> void FlowbndFluxold(Param XParam, double totaltime, BlockP<T>

if (XParam.GPUDEVICE >= 0)
{
//bndFluxGPU << < gridDimBBND, blockDim, 0 >> > (XParam, side, XBlock, Atmp, float(itime), XEv.zs, XEv.h, un, ut, Fh, Fq, S);
//bndFluxGPU <<< gridDimBBND, blockDim, 0 >>> (XParam, side, XBlock, Atmp, float(itime), XEv.zs, XEv.h, un, ut, Fh, Fq, S);
//CUDA_CHECK(cudaDeviceSynchronize());
}
else
Expand Down Expand Up @@ -290,9 +292,9 @@ template <class T> __global__ void bndFluxGPUSide(Param XParam, bndsegmentside s
}





zsbnd = zsbnd + XParam.zsoffset;


int inside = Inside(halowidth, blkmemwidth, side.isright, side.istop, ix, iy, ib);

Expand Down Expand Up @@ -443,6 +445,9 @@ template <class T> void bndFluxGPUSideCPU(Param XParam, bndsegmentside side, Blo
zsbnd = interp2BUQ(XParam.xo + xx, XParam.yo + yy, Zsmap);
}


zsbnd = zsbnd + XParam.zsoffset;


int i = memloc(halowidth, blkmemwidth, ix, iy, ib);
int inside = Inside(halowidth, blkmemwidth, side.isright, side.istop, ix, iy, ib);
Expand Down
24 changes: 12 additions & 12 deletions src/ConserveElevation.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -31,13 +31,13 @@ template <class T> void conserveElevationGPU(Param XParam, BlockP<T> XBlock, Evo
dim3 gridDim(XParam.nblk, 1, 1);


conserveElevationLeft << <gridDim, blockDimHaloLR, 0 >> > (XParam, XBlock, XEv, zb);
conserveElevationLeft <<<gridDim, blockDimHaloLR, 0 >>> (XParam, XBlock, XEv, zb);
CUDA_CHECK(cudaDeviceSynchronize());
conserveElevationRight << <gridDim, blockDimHaloLR, 0 >> > (XParam, XBlock, XEv, zb);
conserveElevationRight <<<gridDim, blockDimHaloLR, 0 >>> (XParam, XBlock, XEv, zb);
CUDA_CHECK(cudaDeviceSynchronize());
conserveElevationTop << <gridDim, blockDimHaloBT, 0 >> > (XParam, XBlock, XEv, zb);
conserveElevationTop <<<gridDim, blockDimHaloBT, 0 >>> (XParam, XBlock, XEv, zb);
CUDA_CHECK(cudaDeviceSynchronize());
conserveElevationBot << <gridDim, blockDimHaloBT, 0 >> > (XParam, XBlock, XEv, zb);
conserveElevationBot <<<gridDim, blockDimHaloBT, 0 >>> (XParam, XBlock, XEv, zb);
CUDA_CHECK(cudaDeviceSynchronize());

}
Expand Down Expand Up @@ -329,13 +329,13 @@ template <class T> void WetDryProlongationGPU(Param XParam, BlockP<T> XBlock, Ev

//WetDryProlongationGPUBot

WetDryProlongationGPULeft << <gridDim, blockDimHaloLR, 0 >> > (XParam, XBlock, XEv, zb);
WetDryProlongationGPULeft <<<gridDim, blockDimHaloLR, 0 >>> (XParam, XBlock, XEv, zb);
CUDA_CHECK(cudaDeviceSynchronize());
WetDryProlongationGPURight << <gridDim, blockDimHaloLR, 0 >> > (XParam, XBlock, XEv, zb);
WetDryProlongationGPURight <<<gridDim, blockDimHaloLR, 0 >>> (XParam, XBlock, XEv, zb);
CUDA_CHECK(cudaDeviceSynchronize());
WetDryProlongationGPUTop << <gridDim, blockDimHaloBT, 0 >> > (XParam, XBlock, XEv, zb);
WetDryProlongationGPUTop <<<gridDim, blockDimHaloBT, 0 >>> (XParam, XBlock, XEv, zb);
CUDA_CHECK(cudaDeviceSynchronize());
WetDryProlongationGPUBot << <gridDim, blockDimHaloBT, 0 >> > (XParam, XBlock, XEv, zb);
WetDryProlongationGPUBot <<<gridDim, blockDimHaloBT, 0 >>> (XParam, XBlock, XEv, zb);
CUDA_CHECK(cudaDeviceSynchronize());

}
Expand All @@ -350,13 +350,13 @@ template <class T> void WetDryRestrictionGPU(Param XParam, BlockP<T> XBlock, Evo

//WetDryProlongationGPUBot

WetDryRestrictionGPULeft << <gridDim, blockDimHaloLR, 0 >> > (XParam, XBlock, XEv, zb);
WetDryRestrictionGPULeft <<<gridDim, blockDimHaloLR, 0 >>> (XParam, XBlock, XEv, zb);
CUDA_CHECK(cudaDeviceSynchronize());
WetDryRestrictionGPURight << <gridDim, blockDimHaloLR, 0 >> > (XParam, XBlock, XEv, zb);
WetDryRestrictionGPURight <<<gridDim, blockDimHaloLR, 0 >>> (XParam, XBlock, XEv, zb);
CUDA_CHECK(cudaDeviceSynchronize());
WetDryRestrictionGPUTop << <gridDim, blockDimHaloBT, 0 >> > (XParam, XBlock, XEv, zb);
WetDryRestrictionGPUTop <<<gridDim, blockDimHaloBT, 0 >>> (XParam, XBlock, XEv, zb);
CUDA_CHECK(cudaDeviceSynchronize());
WetDryRestrictionGPUBot << <gridDim, blockDimHaloBT, 0 >> > (XParam, XBlock, XEv, zb);
WetDryRestrictionGPUBot <<<gridDim, blockDimHaloBT, 0 >>> (XParam, XBlock, XEv, zb);
CUDA_CHECK(cudaDeviceSynchronize());

}
Expand Down
Loading