/* Blocked matrix multiplication.
* TODO slower than transpose, sometimes similar timing to naive.
* Why do they say that this is better?
* */
void mat_mul_cpu_block(const F *A, const F *B, F *C, size_t n, Cache *cache) {
size_t ib, jb, kb, i, j, k, i_max, j_max, k_max, nb ;
F tmp;
UNUSED(cache);
nb = lround(sqrt(n));
for (ib = 0; ib < n; ib += nb) {
i_max = zmin(ib + nb, n);
for (jb = 0; jb < n; jb += nb) {
k_max = zmin(jb + nb, n);
for (kb = 0; kb < n; kb += nb) {
j_max = zmin(kb + nb, n);
/* TODO would be cool to recurse here, but would require more offset parameters,
* likely similar to tose of CBLAS. */
for (i = ib; i < i_max; ++i) {
for (j = kb; j < j_max; ++j) {
tmp = 0.0;
for (k = jb; k < k_max; ++k) {
tmp += A[i*n+k] * B[k*n+j];
}
C[i*n+j] += tmp;
}
}
}
}
}
}
// We create packets and directly fill each zBuffer tile. Note that we
// really store t values
void TaskRayTraceHiZ::run(size_t taskID)
{
const uint32 taskX = taskID % this->taskXNum;
const uint32 taskY = taskID / this->taskXNum;
const uint32 startX = taskX * this->width;
const uint32 startY = taskY * this->height;
const uint32 endX = startX + this->width;
const uint32 endY = startY + this->height;
uint32 tileY = startY / HiZ::Tile::height;
for (uint32 y = startY; y < endY; y += RayPacket::height, ++tileY) {
uint32 tileX = startX / HiZ::Tile::width;
for (uint32 x = startX; x < endX; x += RayPacket::width, ++tileX) {
RayPacket pckt;
PacketHit hit;
gen.generate(pckt, x, y);
intersector->traverse(pckt, hit);
ssef zmin(inf), zmax(neg_inf);
const uint32 tileID = tileX + tileY * zBuffer->tileXNum;
PF_ASSERT(tileID < zBuffer->tileNum);
HiZ::Tile &tile = zBuffer->tiles[tileID];
for (uint32 chunkID = 0; chunkID < HiZ::Tile::chunkNum; ++chunkID) {
//const ssef t = hit.t[chunkID];
const ssef t = hit.t[chunkID] *dot(sse3f(view.x,view.y,view.z), pckt.dir[chunkID]);
tile.z[chunkID] = t;
zmin = min(zmin, t);
zmax = max(zmax, t);
}
tile.zmin = reduce_min(zmin)[0];
tile.zmax = reduce_max(zmax)[0];
}
}
}
/* Like row private, but to reduce global memory accesses,
* we copy only once per work group to local memory.
*
* Each work group contains a few rows of A.
*
* We load the first column, multiply all rows by that column, synrhconize,
* load another column, and so on.
*
* This leads to a thread blockage / memory access tradeoff.
*
* We make work groups as large as possible to reload memory less times. */
void mat_mul_cl_row_priv_col_local(const F *A, const F *B, F *C, size_t n, Cache *cache) {
char options[256];
cl_uint ncl;
size_t global_work_size, local_work_size, mat_sizeof;
/* Setup variables. */
global_work_size = n;
mat_sizeof = n * n * sizeof(F);
ncl = n;
/* Run kernel. */
snprintf(options, sizeof(options), "-DPRIV_ROW_SIZE=%ju", n);
common_create_kernel_file(&cache->common, "matmul_row_priv_col_local.cl", options);
local_work_size = 0;
clGetDeviceInfo(cache->common.device, CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(local_work_size), &local_work_size, NULL);
local_work_size = zmin(local_work_size, n);
clEnqueueWriteBuffer(cache->common.command_queue, cache->buf_a, CL_TRUE, 0, mat_sizeof, (F*)A, 0, NULL, NULL);
clEnqueueWriteBuffer(cache->common.command_queue, cache->buf_b, CL_TRUE, 0, mat_sizeof, (F*)B, 0, NULL, NULL);
clSetKernelArg(cache->common.kernel, 0, sizeof(cache->buf_a), &cache->buf_a);
clSetKernelArg(cache->common.kernel, 1, sizeof(cache->buf_b), &cache->buf_b);
clSetKernelArg(cache->common.kernel, 2, sizeof(cache->buf_c), &cache->buf_c);
clSetKernelArg(cache->common.kernel, 3, n * sizeof(F), NULL);
clSetKernelArg(cache->common.kernel, 4, sizeof(ncl), &ncl);
clEnqueueNDRangeKernel(cache->common.command_queue, cache->common.kernel, 1, NULL, &global_work_size, &local_work_size, 0, NULL, NULL);
clFlush(cache->common.command_queue);
clFinish(cache->common.command_queue);
clEnqueueReadBuffer(cache->common.command_queue, cache->buf_c, CL_TRUE, 0, mat_sizeof, C, 0, NULL, NULL);
}
void Instance::transformBoundingBox()
{
auto b = i->getBoundingBox();
BoundingBox bb;
for(int i = 0; i <= RayTracer::getInstance()->maxTime; ++i)
{
std::set<double> x,y,z;
auto m = makeMatrices(i);
auto p = Vector(b.xmin(i),0,0);
p = transformLoc(m.first, p);
x.insert(p.x);
y.insert(p.y);
z.insert(p.z);
p = Vector(b.xmax(i),0,0);
p = transformLoc(m.first, p);
x.insert(p.x);
y.insert(p.y);
z.insert(p.z);
p = Vector(0,b.ymin(i),0);
p = transformLoc(m.first, p);
x.insert(p.x);
y.insert(p.y);
z.insert(p.z);
p = Vector(0,b.ymax(i),0);
p = transformLoc(m.first, p);
x.insert(p.x);
y.insert(p.y);
z.insert(p.z);
p = Vector(0,0,b.zmin(i));
p = transformLoc(m.first, p);
x.insert(p.x);
y.insert(p.y);
z.insert(p.z);
p = Vector(0,0,b.zmax(i));
p = transformLoc(m.first, p);
x.insert(p.x);
y.insert(p.y);
z.insert(p.z);
bb.xmin.addFrame(i, *x.begin());
bb.xmax.addFrame(i, *x.rbegin());
bb.ymin.addFrame(i, *y.begin());
bb.ymax.addFrame(i, *y.rbegin());
bb.zmin.addFrame(i, *z.begin());
bb.zmax.addFrame(i, *z.rbegin());
}
bbox = bb;
}
/* Copy as many cols from B as possibl to the local memory, only then start multiplying.
* This leads to less memory barrier hits.
* How many rows we copy is limited by the local memory size, ideally the entire matrix will fit. */
void mat_mul_cl_row_priv_cols_local(const F *A, const F *B, F *C, size_t n, Cache *cache) {
char options[256];
cl_uint ncl, n_local_cols;
cl_ulong local_mem_size;
size_t col_size, global_work_size, local_work_size, mat_sizeof;
/* Setup variables. */
col_size = n * sizeof(F);
global_work_size = n;
mat_sizeof = n * n * sizeof(F);
ncl = n;
/* Run kernel. */
snprintf(options, sizeof(options), "-DPRIV_ROW_SIZE=%ju", n);
common_create_kernel_file(&cache->common, "matmul_row_priv_cols_local.cl", options);
local_work_size = 0;
clGetDeviceInfo(cache->common.device, CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(local_work_size), &local_work_size, NULL);
local_work_size = zmin(local_work_size, n);
local_mem_size = 0;
clGetDeviceInfo(cache->common.device, CL_DEVICE_LOCAL_MEM_SIZE, sizeof(local_mem_size), &local_mem_size, NULL);
/* TODO: can blow up without that - 1. Why?
* It only reaches the max without it, not crosses, right?
* So bug in the kernel? */
n_local_cols = zmin(local_mem_size / col_size, n) - 1;
/*puts("");*/
/*printf("max memory %llu\n", (unsigned long long)local_mem_size);*/
/*printf("n_local_cols %llu\n", (unsigned long long)n_local_cols);*/
/*printf("memory %llu\n", (unsigned long long)n_local_cols * n * sizeof(F));*/
clEnqueueWriteBuffer(cache->common.command_queue, cache->buf_a, CL_TRUE, 0, mat_sizeof, (F*)A, 0, NULL, NULL);
clEnqueueWriteBuffer(cache->common.command_queue, cache->buf_b, CL_TRUE, 0, mat_sizeof, (F*)B, 0, NULL, NULL);
clSetKernelArg(cache->common.kernel, 0, sizeof(cache->buf_a), &cache->buf_a);
clSetKernelArg(cache->common.kernel, 1, sizeof(cache->buf_b), &cache->buf_b);
clSetKernelArg(cache->common.kernel, 2, sizeof(cache->buf_c), &cache->buf_c);
clSetKernelArg(cache->common.kernel, 3, n_local_cols * col_size, NULL);
clSetKernelArg(cache->common.kernel, 4, sizeof(ncl), &ncl);
clSetKernelArg(cache->common.kernel, 5, sizeof(n_local_cols), &n_local_cols);
clEnqueueNDRangeKernel(cache->common.command_queue, cache->common.kernel, 1, NULL, &global_work_size, &local_work_size, 0, NULL, NULL);
clFlush(cache->common.command_queue);
clFinish(cache->common.command_queue);
clEnqueueReadBuffer(cache->common.command_queue, cache->buf_c, CL_TRUE, 0, mat_sizeof, C, 0, NULL, NULL);
}
void mat_mul_cl_block(const F *A, const F *B, F *C, size_t n, Cache *cache) {
cl_uint ncl, nblkcl;
size_t global_work_size[2], local_work_size[2], mat_sizeof, nblk;
/* Setup variables. */
global_work_size[0] = n;
global_work_size[1] = n;
mat_sizeof = n * n * sizeof(F);
ncl = n;
/* Run kernel. */
common_create_kernel_file(&cache->common, "matmul_block.cl", NULL);
clGetDeviceInfo(cache->common.device, CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(nblk), &nblk, NULL);
nblk = sqrt(zmin(nblk, n));
nblk = zmin(nblk, 3);
nblkcl = nblk;
local_work_size[0] = nblk;
local_work_size[1] = nblk;
clEnqueueWriteBuffer(cache->common.command_queue, cache->buf_a, CL_TRUE, 0, mat_sizeof, (F*)A, 0, NULL, NULL);
clEnqueueWriteBuffer(cache->common.command_queue, cache->buf_b, CL_TRUE, 0, mat_sizeof, (F*)B, 0, NULL, NULL);
clSetKernelArg(cache->common.kernel, 0, sizeof(cache->buf_a), &cache->buf_a);
clSetKernelArg(cache->common.kernel, 1, sizeof(cache->buf_b), &cache->buf_b);
clSetKernelArg(cache->common.kernel, 2, sizeof(cache->buf_c), &cache->buf_c);
clSetKernelArg(cache->common.kernel, 3, nblk * nblk * sizeof(F), NULL);
printf("nblk = %llu\n", (unsigned long long)nblk);
printf("local memory = %llu\n", (unsigned long long)2 * nblk * nblk * sizeof(F));
clSetKernelArg(cache->common.kernel, 4, nblk * nblk * sizeof(F), NULL);
clSetKernelArg(cache->common.kernel, 5, sizeof(ncl), &ncl);
clSetKernelArg(cache->common.kernel, 6, sizeof(nblkcl), &nblkcl);
clEnqueueNDRangeKernel(
cache->common.command_queue, cache->common.kernel, 2, NULL,
global_work_size, local_work_size, 0, NULL, NULL
);
clFlush(cache->common.command_queue);
clFinish(cache->common.command_queue);
clEnqueueReadBuffer(cache->common.command_queue, cache->buf_c, CL_TRUE, 0, mat_sizeof, C, 0, NULL, NULL);
}
/// Retrieves the map name from 'string' (removing .gat extension if present).
/// Result gets placed either into 'buf' or in a static local buffer.
const char* mapindex_getmapname(const char* string, char* output) {
static char buf[MAP_NAME_LENGTH];
char* dest = (output != NULL) ? output : buf;
size_t len = strnlen(string, MAP_NAME_LENGTH_EXT);
if (len == MAP_NAME_LENGTH_EXT) {
ShowWarning("(mapindex_normalize_name) Map name '%*s' is too long!\n", 2*MAP_NAME_LENGTH_EXT, string);
len--;
}
if (len >= 4 && stricmp(&string[len-4], ".gat") == 0)
len -= 4; // strip .gat extension
len = zmin(len, MAP_NAME_LENGTH-1);
safestrncpy(dest, string, len+1);
memset(&dest[len], '\0', MAP_NAME_LENGTH-len);
return dest;
}
int inter_pet_fromsql(int pet_id, struct s_pet* p)
{
char* data;
size_t len;
#ifdef NOISY
ShowInfo("Loading pet (%d)...\n",pet_id);
#endif
memset(p, 0, sizeof(struct s_pet));
//`pet` (`pet_id`, `class`,`name`,`account_id`,`char_id`,`level`,`egg_id`,`equip`,`intimate`,`hungry`,`rename_flag`,`incubate`)
if( SQL_ERROR == Sql_Query(sql_handle, "SELECT `pet_id`, `class`,`name`,`account_id`,`char_id`,`level`,`egg_id`,`equip`,`intimate`,`hungry`,`rename_flag`,`incubate` FROM `%s` WHERE `pet_id`='%d'", schema_config.pet_db, pet_id) )
{
Sql_ShowDebug(sql_handle);
return 0;
}
if( SQL_SUCCESS == Sql_NextRow(sql_handle) )
{
p->pet_id = pet_id;
Sql_GetData(sql_handle, 1, &data, NULL); p->class_ = atoi(data);
Sql_GetData(sql_handle, 2, &data, &len); memcpy(p->name, data, zmin(len, NAME_LENGTH));
Sql_GetData(sql_handle, 3, &data, NULL); p->account_id = atoi(data);
Sql_GetData(sql_handle, 4, &data, NULL); p->char_id = atoi(data);
Sql_GetData(sql_handle, 5, &data, NULL); p->level = atoi(data);
Sql_GetData(sql_handle, 6, &data, NULL); p->egg_id = atoi(data);
Sql_GetData(sql_handle, 7, &data, NULL); p->equip = atoi(data);
Sql_GetData(sql_handle, 8, &data, NULL); p->intimate = atoi(data);
Sql_GetData(sql_handle, 9, &data, NULL); p->hungry = atoi(data);
Sql_GetData(sql_handle, 10, &data, NULL); p->rename_flag = atoi(data);
Sql_GetData(sql_handle, 11, &data, NULL); p->incubate = atoi(data);
Sql_FreeResult(sql_handle);
p->hungry = cap_value(p->hungry, 0, 100);
p->intimate = cap_value(p->intimate, 0, 1000);
if( charserv_config.save_log )
ShowInfo("Pet loaded (%d - %s).\n", pet_id, p->name);
}
return 0;
}
//------------------------------------------------------------------------------------------------------------------------------------
// called when we want to draw the 3D data in our app.
//------------------------------------------------------------------------------------------------------------------------------------
void draw3D()
{
const float DEG_TO_RAD = PI / 180.0f;
const Vec3 xAxis(1.0f, 0, 0);
const Vec3 yAxis(0, 1.0f, 0);
translate(0, 0, -g_zoom);
translate(g_tx, g_ty, 0);
rotate(g_rotx * DEG_TO_RAD, xAxis);
rotate(g_roty * DEG_TO_RAD, yAxis);
// draw the grid on the floor
setColour(0.25f, 0.25f, 0.25f);
for(float i = -10.0f; i <= 10.1f; i += 1.0f)
{
Vec3 zmin(i, 0, -10);
Vec3 zmax(i, 0, 10);
Vec3 xmin(-10, 0, i);
Vec3 xmax(10, 0, i);
drawLine(xmin, xmax);
drawLine(zmin, zmax);
}
}
int tester()
{
//description
std::vector<std::string> neutrals;
std::vector<std::string> ions;
neutrals.push_back("N2");
neutrals.push_back("CH4");
neutrals.push_back("C2H");
//ionic system contains neutral system
ions = neutrals;
ions.push_back("N2+");
Scalar MN(14.008L), MC(12.011), MH(1.008L);
Scalar MN2 = 2.L*MN , MCH4 = MC + 4.L*MH, MC2H = 2.L * MC + MH;
std::vector<Scalar> Mm;
Mm.push_back(MN2);
Mm.push_back(MCH4);
Mm.push_back(MC2H);
//densities
std::vector<Scalar> molar_frac;
molar_frac.push_back(0.95999L);
molar_frac.push_back(0.04000L);
molar_frac.push_back(0.00001L);
molar_frac.push_back(0.L);
Scalar dens_tot(1e12L);
//hard sphere radius
std::vector<Scalar> hard_sphere_radius;
hard_sphere_radius.push_back(2.0675e-8L * 1e-2L); //N2 in cm -> m
hard_sphere_radius.push_back(2.3482e-8L * 1e-2L); //CH4 in cm -> m
hard_sphere_radius.push_back(0.L); //C2H
//zenith angle
//not necessary
//photon flux
//not necessary
////cross-section
//not necessary
//altitudes
Scalar zmin(600.),zmax(1400.),zstep(10.);
//binary diffusion
Scalar bCN1(1.04e-5 * 1e-4),bCN2(1.76); //cm2 -> m2
Planet::DiffusionType CN_model(Planet::DiffusionType::Wakeham);
Scalar bCC1(5.73e16 * 1e-4),bCC2(0.5); //cm2 -> m2
Planet::DiffusionType CC_model(Planet::DiffusionType::Wilson);
Scalar bNN1(0.1783 * 1e-4),bNN2(1.81); //cm2 -> m2
Planet::DiffusionType NN_model(Planet::DiffusionType::Massman);
/************************
* first level
************************/
//altitude
Planet::Altitude<Scalar,std::vector<Scalar> > altitude(zmin,zmax,zstep);
//neutrals
Antioch::ChemicalMixture<Scalar> neutral_species(neutrals);
//ions
Antioch::ChemicalMixture<Scalar> ionic_species(ions);
//chapman
//not needed
//binary diffusion
Planet::BinaryDiffusion<Scalar> N2N2( Antioch::Species::N2, Antioch::Species::N2 , bNN1, bNN2, NN_model);
Planet::BinaryDiffusion<Scalar> N2CH4( Antioch::Species::N2, Antioch::Species::CH4, bCN1, bCN2, CN_model);
Planet::BinaryDiffusion<Scalar> CH4CH4( Antioch::Species::CH4, Antioch::Species::CH4, bCC1, bCC2, CC_model);
Planet::BinaryDiffusion<Scalar> N2C2H( Antioch::Species::N2, Antioch::Species::C2H);
Planet::BinaryDiffusion<Scalar> CH4C2H( Antioch::Species::CH4, Antioch::Species::C2H);
std::vector<std::vector<Planet::BinaryDiffusion<Scalar> > > bin_diff_coeff;
bin_diff_coeff.resize(2);
bin_diff_coeff[0].push_back(N2N2);
bin_diff_coeff[0].push_back(N2CH4);
bin_diff_coeff[0].push_back(N2C2H);
bin_diff_coeff[1].push_back(N2CH4);
bin_diff_coeff[1].push_back(CH4CH4);
bin_diff_coeff[1].push_back(CH4C2H);
/************************
* second level
************************/
//temperature
std::vector<Scalar> T0,Tz;
read_temperature<Scalar>(T0,Tz,"input/temperature.dat");
std::vector<Scalar> neutral_temperature;
linear_interpolation(T0,Tz,altitude.altitudes(),neutral_temperature);
Planet::AtmosphericTemperature<Scalar, std::vector<Scalar> > temperature(neutral_temperature, neutral_temperature, altitude);
//photon opacity
//not needed
//reaction sets
//not needed
/************************
* third level
************************/
//atmospheric mixture
Planet::AtmosphericMixture<Scalar,std::vector<Scalar>, std::vector<std::vector<Scalar> > > composition(neutral_species, ionic_species, altitude, temperature);
composition.init_composition(molar_frac,dens_tot);
composition.set_hard_sphere_radius(hard_sphere_radius);
composition.initialize();
//kinetics evaluators
//not needed
/************************
* fourth level
************************/
//photon evaluator
//not needed
//molecular diffusion
Planet::MolecularDiffusionEvaluator<Scalar,std::vector<Scalar>, std::vector<std::vector<Scalar> > > molecular_diffusion(bin_diff_coeff,
composition,
altitude,
temperature);
molecular_diffusion.make_molecular_diffusion();
//eddy diffusion
//not needed
/************************
* checks
************************/
molar_frac.pop_back();//get the ion outta here
Scalar Matm(0.L);
for(unsigned int s = 0; s < molar_frac.size(); s++)
{
Matm += molar_frac[s] * composition.neutral_composition().M(s);
}
Matm *= 1e-3L; //to kg
std::vector<std::vector<Scalar> > densities;
calculate_densities(densities, dens_tot, molar_frac, zmin,zmax,zstep, temperature.neutral_temperature(), Mm);
//N2, CH4, C2H
std::vector<std::vector<Scalar> > Dij;
Dij.resize(2);
Dij[0].resize(3,0.L);
Dij[1].resize(3,0.L);
int return_flag(0);
for(unsigned int iz = 0; iz < altitude.altitudes().size(); iz++)
{
Scalar P = pressure(composition.total_density()[iz],temperature.neutral_temperature()[iz]);
Scalar T = temperature.neutral_temperature()[iz];
Dij[0][0] = binary_coefficient(T,P,bNN1,bNN2); //N2 N2
Dij[1][1] = binary_coefficient(T,P,bCC1 * Antioch::ant_pow(Planet::Constants::Convention::T_standard<Scalar>(),bCC2 + Scalar(1.L))
* Planet::Constants::Universal::kb<Scalar>()
/ Planet::Constants::Convention::P_normal<Scalar>(),bCC2 + Scalar(1.L)); //CH4 CH4
Dij[0][1] = binary_coefficient(T,P,bCN1 * Antioch::ant_pow(Planet::Constants::Convention::T_standard<Scalar>(),bCN2),bCN2); //N2 CH4
Dij[0][2] = binary_coefficient(Dij[0][0],Mm[0],Mm[2]); //N2 C2H
Dij[1][2] = binary_coefficient(Dij[1][1],Mm[1],Mm[2]); //CH4 C2H
Dij[1][0] = Dij[0][1]; //CH4 N2
for(unsigned int s = 0; s < molar_frac.size(); s++)
{
Scalar tmp(0.L);
Scalar M_diff(0.L);
for(unsigned int medium = 0; medium < 2; medium++)
{
if(s == medium)continue;
tmp += densities[medium][iz]/Dij[medium][s];
}
Scalar Ds = (barometry(zmin,altitude.altitudes()[iz],neutral_temperature[iz],Matm,dens_tot) - densities[s][iz]) / tmp;
for(unsigned int j = 0; j < molar_frac.size(); j++)
{
if(s == j)continue;
M_diff += composition.total_density()[iz] * composition.neutral_molar_fraction()[j][iz] * composition.neutral_composition().M(j);
}
M_diff /= Scalar(molar_frac.size() - 1);
Scalar Dtilde = Ds / (Scalar(1.L) - composition.neutral_molar_fraction()[s][iz] * (Scalar(1.L) - composition.neutral_composition().M(s)/M_diff));
return_flag = return_flag ||
check_test(Dtilde,molecular_diffusion.Dtilde()[s][iz],"D tilde of species at altitude");
}
return_flag = return_flag ||
check_test(Dij[0][0],molecular_diffusion.binary_coefficient(0,0,T,P),"binary molecular coefficient N2 N2 at altitude") ||
check_test(Dij[0][1],molecular_diffusion.binary_coefficient(0,1,T,P),"binary molecular coefficient N2 CH4 at altitude") ||
check_test(Dij[0][2],molecular_diffusion.binary_coefficient(0,2,T,P),"binary molecular coefficient N2 C2H at altitude") ||
check_test(Dij[1][1],molecular_diffusion.binary_coefficient(1,1,T,P),"binary molecular coefficient CH4 CH4 at altitude") ||
check_test(Dij[1][2],molecular_diffusion.binary_coefficient(1,2,T,P),"binary molecular coefficient CH4 C2H at altitude");
}
return return_flag;
}
/**
* Initializing autotraders from table
*/
void do_init_buyingstore_autotrade( void ) {
if(battle_config.feature_autotrade) {
if (Sql_Query(mmysql_handle,
"SELECT `id`, `account_id`, `char_id`, `sex`, `title`, `limit`, `body_direction`, `head_direction`, `sit` "
"FROM `%s` "
"WHERE `autotrade` = 1 AND `limit` > 0 AND (SELECT COUNT(`buyingstore_id`) FROM `%s` WHERE `buyingstore_id` = `id`) > 0 "
"ORDER BY `id`;",
buyingstores_table, buyingstore_items_table ) != SQL_SUCCESS )
{
Sql_ShowDebug(mmysql_handle);
return;
}
if( Sql_NumRows(mmysql_handle) > 0 ) {
uint16 items = 0;
DBIterator *iter = NULL;
struct s_autotrader *at = NULL;
// Init each autotrader data
while (SQL_SUCCESS == Sql_NextRow(mmysql_handle)) {
size_t len;
char* data;
at = NULL;
CREATE(at, struct s_autotrader, 1);
Sql_GetData(mmysql_handle, 0, &data, NULL); at->id = atoi(data);
Sql_GetData(mmysql_handle, 1, &data, NULL); at->account_id = atoi(data);
Sql_GetData(mmysql_handle, 2, &data, NULL); at->char_id = atoi(data);
Sql_GetData(mmysql_handle, 3, &data, NULL); at->sex = (data[0] == 'F') ? 0 : 1;
Sql_GetData(mmysql_handle, 4, &data, &len); safestrncpy(at->title, data, zmin(len + 1, MESSAGE_SIZE));
Sql_GetData(mmysql_handle, 5, &data, NULL); at->limit = atoi(data);
Sql_GetData(mmysql_handle, 6, &data, NULL); at->dir = atoi(data);
Sql_GetData(mmysql_handle, 7, &data, NULL); at->head_dir = atoi(data);
Sql_GetData(mmysql_handle, 8, &data, NULL); at->sit = atoi(data);
at->count = 0;
if (battle_config.feature_autotrade_direction >= 0)
at->dir = battle_config.feature_autotrade_direction;
if (battle_config.feature_autotrade_head_direction >= 0)
at->head_dir = battle_config.feature_autotrade_head_direction;
if (battle_config.feature_autotrade_sit >= 0)
at->sit = battle_config.feature_autotrade_sit;
// initialize player
CREATE(at->sd, struct map_session_data, 1);
pc_setnewpc(at->sd, at->account_id, at->char_id, 0, gettick(), at->sex, 0);
at->sd->state.autotrade = 1|4;
at->sd->state.monster_ignore = (battle_config.autotrade_monsterignore);
chrif_authreq(at->sd, true);
uidb_put(buyingstore_autotrader_db, at->char_id, at);
}
Sql_FreeResult(mmysql_handle);
// Init items for each autotraders
iter = db_iterator(buyingstore_autotrader_db);
for (at = (struct s_autotrader *)dbi_first(iter); dbi_exists(iter); at = (struct s_autotrader *)dbi_next(iter)) {
uint16 j = 0;
if (SQL_ERROR == Sql_Query(mmysql_handle,
"SELECT `item_id`, `amount`, `price` "
"FROM `%s` "
"WHERE `buyingstore_id` = %d "
"ORDER BY `index` ASC;",
buyingstore_items_table, at->id ) )
{
Sql_ShowDebug(mmysql_handle);
continue;
}
if (!(at->count = (uint16)Sql_NumRows(mmysql_handle))) {
map_quit(at->sd);
buyingstore_autotrader_remove(at, true);
continue;
}
//Init the list
CREATE(at->entries, struct s_autotrade_entry *,at->count);
//Add the item into list
j = 0;
while (SQL_SUCCESS == Sql_NextRow(mmysql_handle) && j < at->count) {
char *data;
CREATE(at->entries[j], struct s_autotrade_entry, 1);
Sql_GetData(mmysql_handle, 0, &data, NULL); at->entries[j]->item_id = atoi(data);
Sql_GetData(mmysql_handle, 1, &data, NULL); at->entries[j]->amount = atoi(data);
Sql_GetData(mmysql_handle, 2, &data, NULL); at->entries[j]->price = atoi(data);
j++;
}
items += j;
Sql_FreeResult(mmysql_handle);
}
dbi_destroy(iter);
ShowStatus("Done loading '"CL_WHITE"%d"CL_RESET"' buyingstore autotraders with '"CL_WHITE"%d"CL_RESET"' items.\n", db_size(buyingstore_autotrader_db), items);
}
}
// Wisp/page request to send
int mapif_parse_WisRequest(int fd)
{
struct WisData* wd;
char name[NAME_LENGTH];
char esc_name[NAME_LENGTH*2+1];// escaped name
char* data;
size_t len;
if ( fd <= 0 ) {return 0;} // check if we have a valid fd
if (RFIFOW(fd,2)-52 >= sizeof(wd->msg)) {
ShowWarning("inter: Wis message size too long.\n");
return 0;
} else if (RFIFOW(fd,2)-52 <= 0) { // normaly, impossible, but who knows...
ShowError("inter: Wis message doesn't exist.\n");
return 0;
}
safestrncpy(name, RFIFOCP(fd,28), NAME_LENGTH); //Received name may be too large and not contain \0! [Skotlex]
Sql_EscapeStringLen(sql_handle, esc_name, name, strnlen(name, NAME_LENGTH));
if( SQL_ERROR == Sql_Query(sql_handle, "SELECT `name` FROM `%s` WHERE `name`='%s'", schema_config.char_db, esc_name) )
Sql_ShowDebug(sql_handle);
// search if character exists before to ask all map-servers
if( SQL_SUCCESS != Sql_NextRow(sql_handle) )
{
unsigned char buf[27];
WBUFW(buf, 0) = 0x3802;
memcpy(WBUFP(buf, 2), RFIFOP(fd, 4), NAME_LENGTH);
WBUFB(buf,26) = 1; // flag: 0: success to send wisper, 1: target character is not loged in?, 2: ignored by target
chmapif_send(fd, buf, 27);
}
else
{// Character exists. So, ask all map-servers
// to be sure of the correct name, rewrite it
Sql_GetData(sql_handle, 0, &data, &len);
memset(name, 0, NAME_LENGTH);
memcpy(name, data, zmin(len, NAME_LENGTH));
// if source is destination, don't ask other servers.
if( strncmp(RFIFOCP(fd,4), name, NAME_LENGTH) == 0 )
{
uint8 buf[27];
WBUFW(buf, 0) = 0x3802;
memcpy(WBUFP(buf, 2), RFIFOP(fd, 4), NAME_LENGTH);
WBUFB(buf,26) = 1; // flag: 0: success to send wisper, 1: target character is not loged in?, 2: ignored by target
chmapif_send(fd, buf, 27);
}
else
{
static int wisid = 0;
CREATE(wd, struct WisData, 1);
// Whether the failure of previous wisp/page transmission (timeout)
check_ttl_wisdata();
wd->id = ++wisid;
wd->fd = fd;
wd->len= RFIFOW(fd,2)-52;
memcpy(wd->src, RFIFOP(fd, 4), NAME_LENGTH);
memcpy(wd->dst, RFIFOP(fd,28), NAME_LENGTH);
memcpy(wd->msg, RFIFOP(fd,52), wd->len);
wd->tick = gettick();
idb_put(wis_db, wd->id, wd);
mapif_wis_message(wd);
}
}
Sql_FreeResult(sql_handle);
return 0;
}
void _Resist_cpu (int idx, int idy, int idz, int idx1, int idy1, int idz1, int idx2, int idy2, int idz2, Field *B1, Field *B2, Field *Emf, Field2D *Eta) {
//<USER_DEFINED>
INPUT(B1);
INPUT(B2);
INPUT(Emf);
INPUT2D(Eta);
OUTPUT(Emf);
//<\USER_DEFINED>
//<EXTERNAL>
real* b1 = B1->field_cpu;
real* b2 = B2->field_cpu;
real* emf = Emf->field_cpu;
real* eta = Eta->field_cpu;
real dx = Dx;
int pitch = Pitch_cpu;
int pitch2d = Pitch2D;
int stride = Stride_cpu;
int size_x = XIP;
int size_y = Ny+2*NGHY-1;
int size_z = Nz+2*NGHZ-1;
//<\EXTERNAL>
//<INTERNAL>
int i;
int j;
int k;
int l1m;
int l2m;
int ll;
real diff1;
real diff2;
//<\INTERNAL>
//<CONSTANT>
// real ymin(Ny+2*NGHY+1);
// real zmin(Nz+2*NGHZ+1);
//<\CONSTANT>
//<MAIN_LOOP>
for (k=1; k<size_z; k++) {
for (j=1; j<size_y; j++) {
for (i=XIM; i<size_x; i++) {
//<#>
ll = l;
l1m = lxm*idx1+lym*idy1+lzm*idz1;
l2m = lxm*idx2+lym*idy2+lzm*idz2;
/* Ideally it should not be the zone size but the distance
between zone centers. The different is very minute here and
does not matter */
diff1 = zone_size_x(j,k)*idx1+zone_size_y(j,k)*idy1+zone_size_z(j,k)*idz1;
diff2 = zone_size_x(j,k)*idx2+zone_size_y(j,k)*idy2+zone_size_z(j,k)*idz2;
emf[ll] += eta[l2D]*((b2[ll]-b2[l1m])/diff1-(b1[ll]-b1[l2m])/diff2);
#ifdef CYLINDRICAL
if (idz1+idz2 == 0) // Considering vertical component of Emf
emf[ll] -= eta[l2D]*(b1[ll]+b1[l2m])*.5/ymin(j);
//Note that (phi,r,z) has a left-handed orientation.
#endif
#ifdef SPHERICAL
if (idy1+idy2 == 0) // Considering radial component of Emf
emf[ll] += eta[l2D]*(b2[ll]+b2[l1m])*.5/ymed(j)*cos(zmin(k))/sin(zmin(k));
if (idz1+idz2 == 0) // Considering colatitude component of Emf
emf[ll] -= eta[l2D]*(b1[ll]+b1[l2m])*.5/ymin(j);
if (idx1+idx2 == 0) // Considering azimuthal component of Emf
emf[ll] += eta[l2D]*(b2[ll]+b2[l1m])*.5/ymin(j);
#endif
//<\#>
}
}
}
//<\MAIN_LOOP>
}
//------------------------------------------------------------------------------------------------------------------------------------
// called when we want to draw the 3D data in our app.
//------------------------------------------------------------------------------------------------------------------------------------
void draw3D()
{
// draw the grid on the floor
setColour(0.25f, 0.25f, 0.25f);
for(float i = -10.0f; i <= 10.1f; i += 1.0f)
{
Vec3 zmin(i, 0, -10);
Vec3 zmax(i, 0, 10);
Vec3 xmin(-10, 0, i);
Vec3 xmax(10, 0, i);
drawLine(xmin, xmax);
drawLine(zmin, zmax);
}
// If using the GPU to compute the lighting (which is what you want to do!)
if(!g_manualLighting)
{
// turn on lighting
enableLighting();
// set the diffuse material colour (this will be modulated by the effect of the lighting)
setColour(1.0f, 1.0f, 1.0f);
// draw the cube geometry
drawPrimitives(g_pointsVN, 24, kQuads);
// turn off lighting
disableLighting();
}
else
{
// otherwise, compute the lighting manually. Don't ever do this in practice! It's insane! (But it may be useful for educational purposes)
// The direction from the vertex to the light (effectively sunlight in this case!).
Vec3 L(-0.6f, 1.0f, -0.2f);
// make sure L has been normalised!
L = normalize(L);
// start drawing some quads
begin(kQuads);
// loop through each vertex normal
for(int i = 0; i < 24; ++i)
{
// compute N.L
// Make sure we clamp this to zero (so that we ignore any negative values).
float N_dot_L = std::max(dot(L, g_pointsVN[i].n), 0.0f);
// the ambient material colour (always gets added to the final colour)
Vec3 Ka(0.2f, 0.2f, 0.2f);
// the diffuse material colour
Vec3 Kd(1.0f, 1.0f, 1.0f);
// Compute the final colour
Vec3 colour = Ka + (Kd * N_dot_L);
// set the vertex colour
setColour(colour);
// specify the vertex
addVertex(g_pointsVN[i].v);
}
// finish drawing our quads
end();
}
// if we are displaying normals
if(g_displayNormals)
{
// make colour pink
setColour(1.0f, 0.0f, 1.0f);
// loop through each vertex
for(int i = 0; i < 24; ++i)
{
// compute an offset (along the normal direction) from the vertex
Vec3 pn = g_pointsVN[i].v + (g_pointsVN[i].n * 0.2f);
// draw a line to show the normal
drawLine(g_pointsVN[i].v, pn);
}
}
}
void StockholmBoundary_cpu(real dt) {
//<USER_DEFINED>
INPUT(Density);
INPUT2D(Density0);
INPUT2D(Density0_avg);
OUTPUT(Density);
#ifdef ADIABATIC
INPUT(Energy);
INPUT2D(Energy0);
OUTPUT(Energy);
#endif
#ifdef X
INPUT(Vx);
INPUT2D(Vx0);
INPUT2D(Vx0_avg);
OUTPUT(Vx);
#endif
#ifdef Y
INPUT(Vy);
INPUT2D(Vy0);
INPUT2D(Vy0_avg);
OUTPUT(Vy);
#endif
#ifdef Z
INPUT(Vz);
INPUT2D(Vz0);
OUTPUT(Vz);
#endif
//<\USER_DEFINED>
//<EXTERNAL>
real* rho = Density->field_cpu;
real* rho0 = Density0->field_cpu;
real* rho0_avg = Density0_avg->field_cpu;
#ifdef X
real* vx = Vx->field_cpu;
real* vx0 = Vx0->field_cpu;
real* vx0_avg = Vx0_avg->field_cpu;
#endif
#ifdef Y
real* vy = Vy->field_cpu;
real* vy0 = Vy0->field_cpu;
real* vy0_avg = Vy0_avg->field_cpu;
#endif
#ifdef Z
real* vz = Vz->field_cpu;
real* vz0 = Vz0->field_cpu;
#endif
#ifdef ADIABATIC
real* e = Energy->field_cpu;
real* e0 = Energy0->field_cpu;
#endif
int pitch = Pitch_cpu;
int stride = Stride_cpu;
int size_x = Nx+2*NGHX;
int size_y = Ny+2*NGHY;
int size_z = Nz+2*NGHZ;
int pitch2d = Pitch2D;
int nghy = NGHY;
real y_min = YMIN;
real y_max = YMAX;
real z_min = ZMIN;
real z_max = ZMAX;
real of = OMEGAFRAME;
real of0 = OMEGAFRAME0;
real r0 = R0;
#ifdef USERWKZ
real wkzin = WKZIN;
real wkzout = WKZOUT;
#else
real wkzin = .0476;
real wkzout = .19;
#endif
int periodic_z = PERIODICZ;
//<\EXTERNAL>
//<INTERNAL>
int i;
int j;
int k;
// Below: we match the standard Stockholm's prescription
// see De Val Borro et al. (2006), section 3.2
// real Y_inf = y_min + (y_max-y_min)*0.0476;
// real Y_sup = y_max - (y_max-y_min)*0.19;
// real Y_inf = exp( log(y_min) + log(y_max/y_min)*0.05);
// real Y_sup = exp( log(y_max) - log(y_max/y_min)*0.10);
real Y_inf = y_min + (y_max-y_min)*wkzin;
real Y_sup = y_max - (y_max-y_min)*wkzout;
real Z_inf = z_min - (z_max-z_min); // Here we push Z_inf & Z_sup
real Z_sup = z_max + (z_max-z_min); // out of the mesh
#ifdef CYLINDRICAL
Z_inf = z_min + (z_max-z_min)*0.1;
Z_sup = z_max - (z_max-z_min)*0.1;
if (periodic_z) { // Push Z_inf & Z_sup out of mesh if periodic in Z
Z_inf = z_min-r0;
Z_sup = z_max+r0;
}
#endif
#ifdef SPHERICAL
Z_inf = M_PI/2.0-(M_PI/2.0-z_min)*1.2;
Z_sup = M_PI/2.0+(M_PI/2.0-z_min)*0.8; // Avoid damping in ghost zones
// if only half upper disk is covered by the mesh
#endif
real radius;
real vx0_target;
real rampy;
real rampz;
real rampzz;
real rampi;
real ramp;
real rho_target;
real vr_target;
real vp_target;
real tau;
real taud;
real ds = 0.03333;
//<\INTERNAL>
//<CONSTANT>
// real xmin(Nx+1);
// real ymin(Ny+2*NGHY+1);
// real zmin(Nz+2*NGHZ+1);
//<\CONSTANT>
//<MAIN_LOOP>
i = j = k = 0;
#ifdef Z
for (k=0; k<size_z; k++) {
#endif
#ifdef Y
for (j=0; j<size_y; j++) {
#endif
#ifdef X
for (i=0; i<size_x; i++) {
#endif
//<#>
rampy = 0.0;
rampz = 0.0;
rampzz = 0.0;
#ifdef Y
if(ymed(j) > Y_sup) {
rampy = (ymed(j)-Y_sup)/(y_max-Y_sup);
#ifdef NOSTOCKHOLMRIGHT
rampy = 0 ;
#endif
#ifdef STOCKHOLMACC
vr_target = vy0_avg[l2D];
vp_target = vx0_avg[l2D];
rho_target = rho0_avg[l2D];
#else
vr_target = vy0[l2D];
vp_target = vx0[l2D];
rho_target = rho0[l2D];
#endif
}
if(ymed(j) < Y_inf) {
rampy = (Y_inf-ymed(j))/(Y_inf-y_min);
vr_target = vy0[l2D];
vp_target = vx0[l2D];
rho_target = rho0[l2D];
/*
#ifdef STOCKHOLMACC
rho_target = rho0_avg[l2D];
#else
rho_target = rho0[l2D];
#endif
*/
}
rampy *= rampy; /* Parabolic ramp as in De Val Borro et al (2006) */
#endif
#ifdef Z
if(zmed(k) > Z_sup) {
rampz = (zmed(k)-Z_sup)/(z_max-Z_sup);
}
if(zmed(k) < Z_inf) {
rampz = (Z_inf-zmed(k))/(Z_inf-z_min);
}
rampz = rampz * rampz; /* vertical ramp in X^2 */
if(zmin(k) > Z_sup) {
rampzz = (zmin(k)-Z_sup)/(z_max-Z_sup);
}
if(zmin(k) < Z_inf) {
rampzz = (Z_inf-zmin(k))/(Z_inf-z_min);
}
rampzz= rampzz * rampzz; /* vertical ramp in X^2 */
#endif
if (periodic_z) {
rampz = 0.0;
rampzz = 0.0;
}
ramp = rampy+rampz;
rampi= rampy+rampzz;
tau = ds*sqrt(ymed(j)*ymed(j)*ymed(j)/G/MSTAR);
if(ramp>0.0) {
taud = tau/ramp;
#ifndef NOWAVEKILLRHO
rho[l] = (rho[l]*taud+rho_target*dt)/(dt+taud);
#endif
#ifndef NOWAVEKILLVPHI
#ifdef X
vx0_target = vx0[l2D];
radius = ymed(j);
#ifdef SPHERICAL
radius *= sin(zmed(k));
#endif
vx0_target -= (of-of0)*radius;
vx[l] = (vx[l]*taud+vx0_target*dt)/(dt+taud);
#endif
#endif
#ifdef Y
vy[l] = (vy[l]*taud+vr_target*dt)/(dt+taud);
#endif
}
#ifdef Z
if(rampi>0.0) {
taud = tau/rampi;
vz[l] = (vz[l]*taud+vz0[l2D]*dt)/(dt+taud);
}
#endif
//<\#>
#ifdef X
}
#endif
#ifdef Y
}
#endif
#ifdef Z
}
#endif
//<\MAIN_LOOP>
}
// Read party from mysql
struct party_data *inter_party_fromsql(int party_id)
{
int leader_id = 0;
int leader_char = 0;
struct party_data* p;
struct party_member* m;
char* data;
size_t len;
int i;
#ifdef NOISY
ShowInfo("Load party request ("CL_BOLD"%d"CL_RESET")\n", party_id);
#endif
if( party_id <= 0 )
return NULL;
//Load from memory
p = (struct party_data*)idb_get(party_db_, party_id);
if( p != NULL )
return p;
p = party_pt;
memset(p, 0, sizeof(struct party_data));
if( SQL_ERROR == Sql_Query(sql_handle, "SELECT `party_id`, `name`,`exp`,`item`, `leader_id`, `leader_char` FROM `%s` WHERE `party_id`='%d'", schema_config.party_db, party_id) )
{
Sql_ShowDebug(sql_handle);
return NULL;
}
if( SQL_SUCCESS != Sql_NextRow(sql_handle) )
return NULL;
p->party.party_id = party_id;
Sql_GetData(sql_handle, 1, &data, &len); memcpy(p->party.name, data, zmin(len, NAME_LENGTH));
Sql_GetData(sql_handle, 2, &data, NULL); p->party.exp = (atoi(data) ? 1 : 0);
Sql_GetData(sql_handle, 3, &data, NULL); p->party.item = atoi(data);
Sql_GetData(sql_handle, 4, &data, NULL); leader_id = atoi(data);
Sql_GetData(sql_handle, 5, &data, NULL); leader_char = atoi(data);
Sql_FreeResult(sql_handle);
// Load members
if( SQL_ERROR == Sql_Query(sql_handle, "SELECT `account_id`,`char_id`,`name`,`base_level`,`last_map`,`online`,`class` FROM `%s` WHERE `party_id`='%d'", schema_config.char_db, party_id) )
{
Sql_ShowDebug(sql_handle);
return NULL;
}
for( i = 0; i < MAX_PARTY && SQL_SUCCESS == Sql_NextRow(sql_handle); ++i )
{
m = &p->party.member[i];
Sql_GetData(sql_handle, 0, &data, NULL); m->account_id = atoi(data);
Sql_GetData(sql_handle, 1, &data, NULL); m->char_id = atoi(data);
Sql_GetData(sql_handle, 2, &data, &len); memcpy(m->name, data, zmin(len, NAME_LENGTH));
Sql_GetData(sql_handle, 3, &data, NULL); m->lv = atoi(data);
Sql_GetData(sql_handle, 4, &data, NULL); m->map = mapindex_name2id(data);
Sql_GetData(sql_handle, 5, &data, NULL); m->online = (atoi(data) ? 1 : 0);
Sql_GetData(sql_handle, 6, &data, NULL); m->class_ = atoi(data);
m->leader = (m->account_id == leader_id && m->char_id == leader_char ? 1 : 0);
}
Sql_FreeResult(sql_handle);
if( charserv_config.save_log )
ShowInfo("Party loaded (%d - %s).\n", party_id, p->party.name);
//Add party to memory.
CREATE(p, struct party_data, 1);
memcpy(p, party_pt, sizeof(struct party_data));
//init state
int_party_calc_state(p);
idb_put(party_db_, party_id, p);
return p;
}
void visctensor_cart_cpu(){
//<USER_DEFINED>
INPUT(Density);
#ifdef X
INPUT(Vx);
OUTPUT(Mmx);
OUTPUT(Mpx);
#endif
#ifdef Y
INPUT(Vy);
OUTPUT(Mmy);
OUTPUT(Mpy);
#endif
#ifdef Z
INPUT(Vz);
OUTPUT(Mmz);
OUTPUT(Mpz);
#endif
//<\USER_DEFINED>
//<EXTERNAL>
real* rho = Density->field_cpu;
#ifdef X
real* vx = Vx->field_cpu;
#endif
#ifdef Y
real* vy = Vy->field_cpu;
#endif
#ifdef Z
real* vz = Vz->field_cpu;
#endif
#ifdef X
real* tauxx = Mmx->field_cpu;
#endif
#ifdef Y
real* tauyy = Mmy->field_cpu;
#endif
#ifdef Z
real* tauzz = Mmz->field_cpu;
#endif
#if defined(X) && defined(Z)
real* tauxz = Mpx->field_cpu;
#endif
#if defined(Y) && defined(X)
real* tauyx = Mpy->field_cpu;
#endif
#if defined(Z) && defined(Y)
real* tauzy = Mpz->field_cpu;
#endif
int pitch = Pitch_cpu;
int stride = Stride_cpu;
int size_x = XIP;
int size_y = Ny+2*NGHY-1;
int size_z = Nz+2*NGHZ-1;
real dx = Dx;
//<\EXTERNAL>
//<INTERNAL>
int i;
int j;
int k;
real div_v;
//<\INTERNAL>
//<CONSTANT>
// real NU(1);
// real Sxj(Ny+2*NGHY);
// real Syj(Ny+2*NGHY);
// real Szj(Ny+2*NGHY);
// real Sxk(Nz+2*NGHZ);
// real Syk(Nz+2*NGHZ);
// real Szk(Nz+2*NGHZ);
// real ymin(Ny+2*NGHY+1);
// real zmin(Nz+2*NGHZ+1);
// real InvVj(Ny+2*NGHY);
//<\CONSTANT>
//<MAIN_LOOP>
i = j = k = 0;
#ifdef Z
for(k=1; k<size_z; k++) {
#endif
#ifdef Y
for(j=1; j<size_y; j++) {
#endif
#ifdef X
for(i=XIM; i<size_x; i++) {
#endif
//<#>
//Evaluate centered divergence.
div_v = 0.0;
#ifdef X
div_v += (vx[lxp]-vx[l])*SurfX(j,k);
#endif
#ifdef Y
div_v += (vy[lyp]*SurfY(j+1,k)-vy[l]*SurfY(j,k));
#endif
#ifdef Z
div_v += (vz[lzp]*SurfZ(j,k+1)-vz[l]*SurfZ(j,k));
#endif
div_v *= 2.0/3.0*InvVol(j,k);
#ifdef X
tauxx[l] = NU*rho[l]*(2.0*(vx[lxp]-vx[l])/dx - div_v);
#endif
#ifdef Y
tauyy[l] = NU*rho[l]*(2.0*(vy[lyp]-vy[l])/(ymin(j+1)-ymin(j)) - div_v);
#endif
#ifdef Z
tauzz[l] = NU*rho[l]*(2.0*(vz[lzp]-vz[l])/(zmin(k+1)-zmin(k)) - div_v);
#endif
#if defined(X) && defined(Z)
tauxz[l] = NU*.25*(rho[l]+rho[lzm]+rho[lxm]+rho[lxm-stride])*((vx[l]-vx[lzm])/(zmed(k)-zmed(k-1)) + (vz[l]-vz[lxm])/dx); //centered on lower, left "radial" edge in y
#endif
#if defined(Y) && defined(X)
tauyx[l] = NU*.25*(rho[l]+rho[lxm]+rho[lym]+rho[lxm-pitch])*((vy[l]-vy[lxm])/dx + (vx[l]-vx[lym])/(ymed(j)-ymed(j-1))); //centered on left, inner vertical edge in z
#endif
#if defined(Z) && defined(Y)
tauzy[l] = NU*.25*(rho[l]+rho[lym]+rho[lzm]+rho[lym-stride])*((vz[l]-vz[lym])/(ymed(j)-ymed(j-1)) + (vy[l]-vy[lzm])/(zmed(k)-zmed(k-1))); //centered on lower, inner edge in x ("azimuthal")
#endif
//<\#>
#ifdef X
}
#endif
#ifdef Y
}
#endif
#ifdef Z
}
#endif
//<\MAIN_LOOP>
}
