Skip to content

Commit

Permalink
fin
Browse files Browse the repository at this point in the history
  • Loading branch information
archibate committed Feb 4, 2022
1 parent ff3f532 commit f63c8cd
Show file tree
Hide file tree
Showing 2 changed files with 6 additions and 4 deletions.
10 changes: 6 additions & 4 deletions 09/01_texture/08/main.cu
Original file line number Diff line number Diff line change
Expand Up @@ -108,7 +108,7 @@ __global__ void sumloss_kernel(CudaSurfaceAccessor<float> sufDiv, float *sum, un
sufPreNext.write(preNext, x, y, z);
}*/

__global__ void heatup_kernel(CudaSurfaceAccessor<float4> sufVel, CudaSurfaceAccessor<float> sufTmp, CudaSurfaceAccessor<float> sufClr, CudaSurfaceAccessor<char> sufBound, float heatRate, float clrRate, unsigned int n) {
__global__ void heatup_kernel(CudaSurfaceAccessor<float4> sufVel, CudaSurfaceAccessor<float> sufTmp, CudaSurfaceAccessor<float> sufClr, CudaSurfaceAccessor<char> sufBound, float tmpAmbient, float heatRate, float clrRate, unsigned int n) {
int x = threadIdx.x + blockDim.x * blockIdx.x;
int y = threadIdx.y + blockDim.y * blockIdx.y;
int z = threadIdx.z + blockDim.z * blockIdx.z;
Expand All @@ -118,7 +118,7 @@ __global__ void heatup_kernel(CudaSurfaceAccessor<float4> sufVel, CudaSurfaceAcc
float4 vel = sufVel.read(x, y, z);
float tmp = sufTmp.read(x, y, z);
float clr = sufClr.read(x, y, z);
vel.z += heatRate * tmp;
vel.z += heatRate * (tmp - tmpAmbient);
vel.z -= clrRate * clr;
sufVel.write(vel, x, y, z);
}
Expand Down Expand Up @@ -300,7 +300,7 @@ struct SmokeSim : DisableCopy {
}

void projection() {
heatup_kernel<<<dim3((n + 7) / 8, (n + 7) / 8, (n + 7) / 8), dim3(8, 8, 8)>>>(vel->accessSurface(), tmp->accessSurface(), clr->accessSurface(), bound->accessSurface(), 0.018f, 0.004f, n);
heatup_kernel<<<dim3((n + 7) / 8, (n + 7) / 8, (n + 7) / 8), dim3(8, 8, 8)>>>(vel->accessSurface(), tmp->accessSurface(), clr->accessSurface(), bound->accessSurface(), 0.05f, 0.018f, 0.004f, n);
divergence_kernel<<<dim3((n + 7) / 8, (n + 7) / 8, (n + 7) / 8), dim3(8, 8, 8)>>>(vel->accessSurface(), div->accessSurface(), bound->accessSurface(), n);
vcycle(0, pre.get(), div.get());
subgradient_kernel<<<dim3((n + 7) / 8, (n + 7) / 8, (n + 7) / 8), dim3(8, 8, 8)>>>(pre->accessSurface(), vel->accessSurface(), bound->accessSurface(), n);
Expand All @@ -316,8 +316,10 @@ struct SmokeSim : DisableCopy {
std::swap(clr, clrNext);
std::swap(tmp, tmpNext);

decay_kernel<<<dim3((n + 7) / 8, (n + 7) / 8, (n + 7) / 8), dim3(8, 8, 8)>>>(tmp->accessSurface(), tmpNext->accessSurface(), bound->accessSurface(), std::exp(-0.1f), std::exp(-0.01f), n);
decay_kernel<<<dim3((n + 7) / 8, (n + 7) / 8, (n + 7) / 8), dim3(8, 8, 8)>>>(tmp->accessSurface(), tmpNext->accessSurface(), bound->accessSurface(), std::exp(-0.5f), std::exp(-0.0003f), n);
decay_kernel<<<dim3((n + 7) / 8, (n + 7) / 8, (n + 7) / 8), dim3(8, 8, 8)>>>(clr->accessSurface(), clrNext->accessSurface(), bound->accessSurface(), std::exp(-0.05f), std::exp(-0.003f), n);
std::swap(tmp, tmpNext);
std::swap(clr, clrNext);
}

void step(int times = 16) {
Expand Down
Binary file modified 09/slides.pptx
Binary file not shown.

0 comments on commit f63c8cd

Please sign in to comment.