diff --git a/examples_ptx/aobench/Makefile_gpu b/examples_ptx/aobench/Makefile_gpu index 5d21a06f..09f62c55 100644 --- a/examples_ptx/aobench/Makefile_gpu +++ b/examples_ptx/aobench/Makefile_gpu @@ -4,7 +4,7 @@ CU_SRC=ao.cu CXX_SRC=ao.cpp ao_serial.cpp PTXCC_REGMAX=64 -LLVM_GPU=1 +# LLVM_GPU=1 NVVM_GPU=1 include ../common_gpu.mk diff --git a/examples_ptx/mandelbrot_tasks/mandelbrot_tasks.ispc b/examples_ptx/mandelbrot_tasks/mandelbrot_tasks.ispc index 17b7b7bd..9defbc5b 100644 --- a/examples_ptx/mandelbrot_tasks/mandelbrot_tasks.ispc +++ b/examples_ptx/mandelbrot_tasks/mandelbrot_tasks.ispc @@ -82,8 +82,8 @@ mandelbrot_ispc(uniform float x0, uniform float y0, uniform float x1, uniform float y1, uniform int width, uniform int height, uniform int maxIterations, uniform int output[]) { - uniform float dx = (x1 - x0) / width; - uniform float dy = (y1 - y0) / height; + const uniform float dx = (x1 - x0) / width; + const uniform float dy = (y1 - y0) / height; const uniform int xspan = max(32, programCount*2); /* make sure it is big enough to avoid false-sharing */ const uniform int yspan = 16; diff --git a/examples_ptx/rt/rt.ispc b/examples_ptx/rt/rt.ispc index 80ce5ff6..4730554d 100644 --- a/examples_ptx/rt/rt.ispc +++ b/examples_ptx/rt/rt.ispc @@ -38,6 +38,13 @@ typedef bool bool_t; #endif typedef float<3> float3; +#ifdef __NVPTX__ +#define uniform_t varying +#else +#define uniform_t uniform +#endif + + struct int3 { @@ -124,8 +131,8 @@ inline #endif static bool_t BBoxIntersect(const uniform float bounds[2][3], const Ray &ray) { - uniform float3 bounds0 = { bounds[0][0], bounds[0][1], bounds[0][2] }; - uniform float3 bounds1 = { bounds[1][0], bounds[1][1], bounds[1][2] }; + const uniform float3 bounds0 = { bounds[0][0], bounds[0][1], bounds[0][2] }; + const uniform float3 bounds1 = { bounds[1][0], bounds[1][1], bounds[1][2] }; float t0 = ray.mint, t1 = ray.maxt; // Check all three axis-aligned slabs. Don't try to early out; it's @@ -164,12 +171,12 @@ static bool_t BBoxIntersect(const uniform float bounds[2][3], #if 1 inline #endif -static bool_t TriIntersect(const uniform Triangle &tri, Ray &ray) { - uniform float3 p0 = { tri.p[0][0], tri.p[0][1], tri.p[0][2] }; - uniform float3 p1 = { tri.p[1][0], tri.p[1][1], tri.p[1][2] }; - uniform float3 p2 = { tri.p[2][0], tri.p[2][1], tri.p[2][2] }; - uniform float3 e1 = p1 - p0; - uniform float3 e2 = p2 - p0; +static bool_t TriIntersect(const uniform_t Triangle tri, Ray &ray) { + const uniform_t float3 p0 = { tri.p[0][0], tri.p[0][1], tri.p[0][2] }; + const uniform_t float3 p1 = { tri.p[1][0], tri.p[1][1], tri.p[1][2] }; + const uniform_t float3 p2 = { tri.p[2][0], tri.p[2][1], tri.p[2][2] }; + const uniform_t float3 e1 = p1 - p0; + const uniform_t float3 e2 = p2 - p0; float3 s1 = Cross(ray.dir, e2); float divisor = Dot(s1, e1); @@ -218,13 +225,13 @@ BVHIntersect(const uniform LinearBVHNode nodes[], while (true) { // Check ray against BVH node - uniform LinearBVHNode node = nodes[nodeNum]; + const uniform LinearBVHNode node = nodes[nodeNum]; if (any(BBoxIntersect(node.bounds, ray))) { - uniform unsigned int nPrimitives = node.nPrimitives; + const uniform unsigned int nPrimitives = node.nPrimitives; if (nPrimitives > 0) { // Intersect ray with primitives in leaf BVH node - uniform unsigned int primitivesOffset = node.offset; - for (uniform unsigned int i = 0; i < nPrimitives; ++i) { + const uniform unsigned int primitivesOffset = node.offset; + for (uniform_t unsigned int i = 0; i < nPrimitives; ++i) { if (TriIntersect(tris[primitivesOffset+i], ray)) hit = true; } @@ -277,8 +284,8 @@ static void raytrace_tile(uniform int x0, uniform int x1, uniform float image[], uniform int id[], const uniform LinearBVHNode nodes[], const uniform Triangle triangles[]) { - uniform float widthScale = (float)(baseWidth) / (float)(width); - uniform float heightScale = (float)(baseHeight) / (float)(height); + const uniform float widthScale = (float)(baseWidth) / (float)(width); + const uniform float heightScale = (float)(baseHeight) / (float)(height); foreach_tiled (y = y0 ... y1, x = x0 ... x1) { Ray ray; @@ -313,12 +320,12 @@ task void raytrace_tile_task(uniform int width, uniform int height, uniform float image[], uniform int id[], const uniform LinearBVHNode nodes[], const uniform Triangle triangles[]) { - uniform int dx = 64, dy = 8; // must match dx, dy below - uniform int xBuckets = (width + (dx-1)) / dx; - uniform int x0 = (taskIndex % xBuckets) * dx; - uniform int x1 = min(x0 + dx, width); - uniform int y0 = (taskIndex / xBuckets) * dy; - uniform int y1 = min(y0 + dy, height); + const uniform int dx = 64, dy = 8; // must match dx, dy below + const uniform int xBuckets = (width + (dx-1)) / dx; + const uniform int x0 = (taskIndex % xBuckets) * dx; + const uniform int x1 = min(x0 + dx, width); + const uniform int y0 = (taskIndex / xBuckets) * dy; + const uniform int y1 = min(y0 + dy, height); raytrace_tile(x0, x1, y0, y1, width, height, baseWidth, baseHeight, raster2camera, camera2world, image, @@ -333,10 +340,10 @@ export void raytrace_ispc_tasks(uniform int width, uniform int height, uniform float image[], uniform int id[], const uniform LinearBVHNode nodes[], const uniform Triangle triangles[]) { - uniform int dx = 64, dy = 8; - uniform int xBuckets = (width + (dx-1)) / dx; - uniform int yBuckets = (height + (dy-1)) / dy; - uniform int nTasks = xBuckets * yBuckets; + const uniform int dx = 64, dy = 8; + const uniform int xBuckets = (width + (dx-1)) / dx; + const uniform int yBuckets = (height + (dy-1)) / dy; + const uniform int nTasks = xBuckets * yBuckets; launch[nTasks] raytrace_tile_task(width, height, baseWidth, baseHeight, raster2camera, camera2world, image, id, nodes, triangles); diff --git a/examples_ptx/volume_rendering/volume.ispc b/examples_ptx/volume_rendering/volume.ispc index 6d0bd8f5..68ba2dfe 100644 --- a/examples_ptx/volume_rendering/volume.ispc +++ b/examples_ptx/volume_rendering/volume.ispc @@ -181,7 +181,7 @@ transmittance(uniform float3 p0, float3 p1, uniform float3 pMin, float tau = 0; float rayLength = sqrt(ray.dir.x * ray.dir.x + ray.dir.y * ray.dir.y + ray.dir.z * ray.dir.z); - uniform float stepDist = 0.2; + const uniform float stepDist = 0.2; float stepT = stepDist / rayLength; float t = rayT0; @@ -207,8 +207,8 @@ distanceSquared(float3 a, float3 b) { static inline float raymarch(uniform float density[], uniform int nVoxels[3], Ray ray) { float rayT0, rayT1; - uniform float3 pMin = {.3, -.2, .3}, pMax = {1.8, 2.3, 1.8}; - uniform float3 lightPos = { -1, 4, 1.5 }; + const uniform float3 pMin = {.3, -.2, .3}, pMax = {1.8, 2.3, 1.8}; + const uniform float3 lightPos = { -1, 4, 1.5 }; if (!IntersectP(ray, pMin, pMax, rayT0, rayT1)) return 0.; @@ -217,11 +217,11 @@ raymarch(uniform float density[], uniform int nVoxels[3], Ray ray) { // Parameters that define the volume scattering characteristics and // sampling rate for raymarching - uniform float Le = .25; // Emission coefficient - uniform float sigma_a = 10; // Absorption coefficient - uniform float sigma_s = 10; // Scattering coefficient - uniform float stepDist = 0.025; // Ray step amount - uniform float lightIntensity = 40; // Light source intensity + const uniform float Le = .25; // Emission coefficient + const uniform float sigma_a = 10; // Absorption coefficient + const uniform float sigma_s = 10; // Scattering coefficient + const uniform float stepDist = 0.025; // Ray step amount + const uniform float lightIntensity = 40; // Light source intensity float tau = 0.f; // accumulated beam transmittance float L = 0; // radiance along the ray @@ -375,15 +375,14 @@ volume_task(uniform float density[], uniform int _nVoxels[3], #define camera2world _camera2world #endif - uniform int dx = 8, dy = 8; // must match value in volume_ispc_tasks - uniform int xbuckets = (width + (dx-1)) / dx; - uniform int ybuckets = (height + (dy-1)) / dy; + const uniform int dx = 8, dy = 8; // must match value in volume_ispc_tasks + const uniform int xbuckets = (width + (dx-1)) / dx; + const uniform int ybuckets = (height + (dy-1)) / dy; - uniform int x0 = (taskIndex % xbuckets) * dx; - uniform int y0 = (taskIndex / xbuckets) * dy; - uniform int x1 = x0 + dx, y1 = y0 + dy; - x1 = min(x1, width); - y1 = min(y1, height); + const uniform int x0 = (taskIndex % xbuckets) * dx; + const uniform int y0 = (taskIndex / xbuckets) * dy; + const uniform int x1 = min(x0 + dx, width); + const uniform int y1 = min(y0 + dy, height); volume_tile(x0, y0, x1, y1, density, nVoxels, raster2camera, camera2world, width, height, image); @@ -406,8 +405,8 @@ volume_ispc_tasks(uniform float density[], uniform int nVoxels[3], const uniform float camera2world[4][4], uniform int width, uniform int height, uniform float image[]) { // Launch tasks to work on (dx,dy)-sized tiles of the image - uniform int dx = 8, dy = 8; - uniform int nTasks = ((width+(dx-1))/dx) * ((height+(dy-1))/dy); + const uniform int dx = 8, dy = 8; + const uniform int nTasks = ((width+(dx-1))/dx) * ((height+(dy-1))/dy); launch[nTasks] volume_task(density, nVoxels, raster2camera, camera2world, width, height, image); sync; diff --git a/stmt.cpp b/stmt.cpp index c1b063bc..1caca0e1 100644 --- a/stmt.cpp +++ b/stmt.cpp @@ -152,19 +152,6 @@ static llvm::Value* lConvertToGenericPtr(FunctionEmitContext *ctx, llvm::Value * return value; llvm::Type *elTy = pt->getElementType(); - assert(elTy->isArrayTy()); - const int numElTot = elTy->getArrayNumElements(); - const int numEl = numElTot/4; -#if 0 - fprintf(stderr, " --- detected addrspace(3) sz= %d --- \n", numEl); -#endif - llvm::ArrayType *arrTy = llvm::dyn_cast(pt->getArrayElementType()); - assert(arrTy != NULL); - llvm::Type *arrElTy = arrTy->getElementType(); -#if 0 - if (arrElTy->isArrayTy()) - Error(currentPos, "Currently \"nvptx\" target doesn't support array-of-array"); -#endif /* convert elTy addrspace(3)* to i64* addrspace(3)* */ llvm::PointerType *Int64Ptr3 = llvm::PointerType::get(LLVMTypes::Int64Type, addressSpace); @@ -177,13 +164,27 @@ static llvm::Value* lConvertToGenericPtr(FunctionEmitContext *ctx, llvm::Value * __cvt2gen_args.push_back(value); value = llvm::CallInst::Create(__cvt2gen, __cvt2gen_args, "gep2gen_cvt", ctx->GetCurrentBasicBlock()); - /* convert i64* to errElTy* */ - llvm::PointerType *arrElTyPt0 = llvm::PointerType::get(arrElTy, 0); - value = ctx->BitCastInst(value, arrElTyPt0, "gep2gen_cast2"); - /* compute offset */ if (addressSpace == 3) { + assert(elTy->isArrayTy()); + const int numElTot = elTy->getArrayNumElements(); + const int numEl = numElTot/4; +#if 0 + fprintf(stderr, " --- detected addrspace(3) sz= %d --- \n", numEl); +#endif + llvm::ArrayType *arrTy = llvm::dyn_cast(pt->getArrayElementType()); + assert(arrTy != NULL); + llvm::Type *arrElTy = arrTy->getElementType(); +#if 0 + if (arrElTy->isArrayTy()) + Error(currentPos, "Currently \"nvptx\" target doesn't support array-of-array"); +#endif + + /* convert i64* to errElTy* */ + llvm::PointerType *arrElTyPt0 = llvm::PointerType::get(arrElTy, 0); + value = ctx->BitCastInst(value, arrElTyPt0, "gep2gen_cast2"); + llvm::Function *funcTid = m->module->getFunction("__tid_x"); llvm::Function *funcWarpSz = m->module->getFunction("__warpsize"); llvm::Value *tid = ctx->CallInst(funcTid, NULL, std::vector(), "gep2gen_tid"); @@ -268,12 +269,14 @@ DeclStmt::EmitCode(FunctionEmitContext *ctx) const { Error(sym->pos, "Non-constant static variable ""\"%s\" is not supported with ""\"nvptx\" target.", sym->name.c_str()); -#if 0 if (g->target->getISA() == Target::NVPTX && sym->type->IsVaryingType()) - Error(sym->pos, - "const static varying variable ""\"%s\" is not supported with ""\"nvptx\" target.", + PerformanceWarning(sym->pos, + "\"const static varying\" variable ""\"%s\" is stored in __global address space with ""\"nvptx\" target.", + sym->name.c_str()); + if (g->target->getISA() == Target::NVPTX && sym->type->IsUniformType()) + PerformanceWarning(sym->pos, + "\"const static uniform\" variable ""\"%s\" is stored in __constant address space with ""\"nvptx\" target.", sym->name.c_str()); -#endif // For static variables, we need a compile-time constant value // for its initializer; if there's no initializer, we use a @@ -310,208 +313,78 @@ DeclStmt::EmitCode(FunctionEmitContext *ctx) const { llvm::GlobalValue::InternalLinkage, cinit, llvm::Twine("static_") + llvm::Twine(sym->pos.first_line) + - llvm::Twine("_") + sym->name.c_str()); + llvm::Twine("_") + sym->name.c_str(), + NULL, + llvm::GlobalVariable::NotThreadLocal, + /*AddressSpace=*/ sym->type->IsUniformType() ? 4 : 0); + sym->storagePtr = lConvertToGenericPtr(ctx, sym->storagePtr, sym->pos); // Tell the FunctionEmitContext about the variable ctx->EmitVariableDebugInfo(sym); } - else { -#if 0 - if (sym->type->IsUniformType() && - sym->type->IsArrayType() && - g->target->getISA() == Target::NVPTX) + else if (sym->type->IsUniformType() && + /* NVPTX: + * only non-constant uniform data types are stored in shared memory + * constant uniform are automatically promoted to varying + */ + !sym->type->IsConstType() && +#if 1 + sym->type->IsArrayType() && +#endif + g->target->getISA() == Target::NVPTX) { - /* deal with "const uniform" or "uniform" arrays for nvptx target */ - if (!sym->type->IsConstType()) - PerformanceWarning(sym->pos, - "\"uniform\" arrays might be slow with \"nvptx\" target. " - "Unless data sharing between program instances is required, use \"varying\" or \"uniform new\"+\"delete\" if possible."); - if (initExpr != NULL && !sym->type->IsConstType()) - Error(initExpr->pos, - "It is not possible to initialize non-constant \"uniform\" array \"%s\" with \"nvptx\" target. " - "Please use \"varying\", \"const static uniform\" or define initializer in the global scope.", - sym->name.c_str()); + PerformanceWarning(sym->pos, + "Non-constant \"uniform\" data types might be slow with \"nvptx\" target. " + "Unless data sharing between program instances is desired, try \"const [static] uniform\", \"varying\" or \"uniform new uniform \"+\"delete\" if possible."); - if (sym->type->IsConstType()) + /* with __shared__ memory everything must be an array */ + int nel = 4; + ArrayType *nat; + if (sym->type->IsArrayType()) { - llvm::Constant *cinit = NULL; - if (initExpr != NULL) { - if (PossiblyResolveFunctionOverloads(initExpr, sym->type) == false) - continue; - // FIXME: we only need this for function pointers; it was - // already done for atomic types and enums in - // DeclStmt::TypeCheck()... - if (dynamic_cast(initExpr) == NULL) { - initExpr = TypeConvertExpr(initExpr, sym->type, - "initializer"); - // FIXME: and this is only needed to re-establish - // constant-ness so that GetConstant below works for - // constant artithmetic expressions... - initExpr = ::Optimize(initExpr); - } - - cinit = initExpr->GetConstant(sym->type); - if (cinit == NULL) - Error(initExpr->pos, "Initializer for static variable " - "\"%s\" must be a constant.", sym->name.c_str()); - } - else - { - Error(sym->pos, "Missing initializer for const variable " - "\"%s\".", sym->name.c_str()); - } - if (cinit == NULL) - cinit = llvm::Constant::getNullValue(llvmType); - - sym->storagePtr = - new llvm::GlobalVariable(*m->module, llvmType, - sym->type->IsConstType(), - llvm::GlobalValue::InternalLinkage, - cinit, - llvm::Twine("local_") + - llvm::Twine(sym->pos.first_line) + - llvm::Twine("_") + sym->name.c_str(), - NULL, - llvm::GlobalVariable::NotThreadLocal, - /*AddressSpace=*/4); /* constant address space */ - // Tell the FunctionEmitContext about the variable - ctx->EmitVariableDebugInfo(sym); - } - else - { - /* fails if pointer passed to function argument, need conversion beforehand */ - llvm::Constant *cinit = NULL; const ArrayType *at = CastType(sym->type); - const int nel = at->GetElementCount(); + nel = at->GetElementCount(); /* we must scale # elements by 4, because a thread-block will run 4 warps * or 128 threads. * ***note-to-me***:please define these value (128threads/4warps) * in nvptx-target definition * instead of compile-time constants */ - const int nel4 = nel*4; - ArrayType nat(at->GetElementType(), nel4); - llvm::Type *llvmType = nat.LLVMType(g->ctx); - cinit = llvm::UndefValue::get(llvmType); - - sym->storagePtr = - new llvm::GlobalVariable(*m->module, llvmType, - sym->type->IsConstType(), - llvm::GlobalValue::InternalLinkage, - cinit, - llvm::Twine("local_") + - llvm::Twine(sym->pos.first_line) + - llvm::Twine("_") + sym->name.c_str(), - NULL, - llvm::GlobalVariable::NotThreadLocal, - /*AddressSpace=*/3); - // Tell the FunctionEmitContext about the variable - ctx->EmitVariableDebugInfo(sym); + nel *= 4; + nat = new ArrayType(at->GetElementType(), nel); } + else + nat = new ArrayType(sym->type, nel); + + llvm::Type *llvmTypeUn = nat->LLVMType(g->ctx); + llvm::Constant *cinit = llvm::UndefValue::get(llvmTypeUn); + + sym->storagePtr = + new llvm::GlobalVariable(*m->module, llvmTypeUn, + sym->type->IsConstType(), + llvm::GlobalValue::InternalLinkage, + cinit, + llvm::Twine("local_") + + llvm::Twine(sym->pos.first_line) + + llvm::Twine("_") + sym->name.c_str(), + NULL, + llvm::GlobalVariable::NotThreadLocal, + /*AddressSpace=*/3); + sym->storagePtr = lConvertToGenericPtr(ctx, sym->storagePtr, sym->pos); + llvm::PointerType *ptrTy = llvm::PointerType::get(sym->type->LLVMType(g->ctx),0); + sym->storagePtr = ctx->BitCastInst(sym->storagePtr, ptrTy, "uniform_decl"); + + // Tell the FunctionEmitContext about the variable; must do + // this before the initializer stuff. + ctx->EmitVariableDebugInfo(sym); + + if (initExpr == 0 && sym->type->IsConstType()) + Error(sym->pos, "Missing initializer for const variable " + "\"%s\".", sym->name.c_str()); + + // And then get it initialized... + sym->parentFunction = ctx->GetFunction(); + InitSymbol(sym->storagePtr, sym->type, initExpr, ctx, sym->pos); } - else if ( - sym->type->IsUniformType() && - g->target->getISA() == Target::NVPTX) - { -#if 1 - // For non-static variables, allocate storage on the stack - sym->storagePtr = ctx->AllocaInst(llvmType, sym->name.c_str()); -#else - PerformanceWarning(sym->pos, - "\"uniform\" variables might be slow with \"nvptx\" target. " - "Please use \"varying\" if possible."); - - ArrayType nat(sym->type, 4); - llvm::Type *llvmType = nat.LLVMType(g->ctx); - llvm::Constant *cinit = llvm::UndefValue::get(llvmType); - - sym->storagePtr = - new llvm::GlobalVariable(*m->module, llvmType, - sym->type->IsConstType(), - llvm::GlobalValue::InternalLinkage, - cinit, - llvm::Twine("local_") + - llvm::Twine(sym->pos.first_line) + - llvm::Twine("_") + sym->name.c_str(), - NULL, - llvm::GlobalVariable::NotThreadLocal, - /*AddressSpace=*/3); - sym->storagePtr = lConvertToGenericPtr(ctx, sym->storagePtr, sym->pos); - llvm::PointerType *ptrTy = - llvm::PointerType::get(sym->type->LLVMType(g->ctx),0); - sym->storagePtr = ctx->BitCastInst(sym->storagePtr, ptrTy, "uniform_alloc"); -#endif - - - // Tell the FunctionEmitContext about the variable; must do - // this before the initializer stuff. - ctx->EmitVariableDebugInfo(sym); - - if (initExpr == 0 && sym->type->IsConstType()) - Error(sym->pos, "Missing initializer for const variable " - "\"%s\".", sym->name.c_str()); - - // And then get it initialized... - sym->parentFunction = ctx->GetFunction(); - InitSymbol(sym->storagePtr, sym->type, initExpr, ctx, sym->pos); - } -#else - if (sym->type->IsUniformType() && - sym->type->IsArrayType() && - g->target->getISA() == Target::NVPTX) - { - PerformanceWarning(sym->pos, - "\"uniform\" data types might be slow with \"nvptx\" target. " - "Unless data sharing between program instances is desired, try either \"varying\" or \"uniform new uniform \"+\"delete\" if possible."); - - int nel = 4; - ArrayType *nat; - if (sym->type->IsArrayType()) - { - const ArrayType *at = CastType(sym->type); - nel = at->GetElementCount(); - /* we must scale # elements by 4, because a thread-block will run 4 warps - * or 128 threads. - * ***note-to-me***:please define these value (128threads/4warps) - * in nvptx-target definition - * instead of compile-time constants - */ - nel *= 4; - nat = new ArrayType(at->GetElementType(), nel); - } - else - nat = new ArrayType(sym->type, nel); - llvm::Type *llvmType = nat->LLVMType(g->ctx); - llvm::Constant *cinit = llvm::UndefValue::get(llvmType); - - sym->storagePtr = - new llvm::GlobalVariable(*m->module, llvmType, - sym->type->IsConstType(), - llvm::GlobalValue::InternalLinkage, - cinit, - llvm::Twine("local_") + - llvm::Twine(sym->pos.first_line) + - llvm::Twine("_") + sym->name.c_str(), - NULL, - llvm::GlobalVariable::NotThreadLocal, - /*AddressSpace=*/3); - sym->storagePtr = lConvertToGenericPtr(ctx, sym->storagePtr, sym->pos); - llvm::PointerType *ptrTy = - llvm::PointerType::get(sym->type->LLVMType(g->ctx),0); - sym->storagePtr = ctx->BitCastInst(sym->storagePtr, ptrTy, "uniform_alloc"); - - // Tell the FunctionEmitContext about the variable; must do - // this before the initializer stuff. - ctx->EmitVariableDebugInfo(sym); - - if (initExpr == 0 && sym->type->IsConstType()) - Error(sym->pos, "Missing initializer for const variable " - "\"%s\".", sym->name.c_str()); - - // And then get it initialized... - sym->parentFunction = ctx->GetFunction(); - InitSymbol(sym->storagePtr, sym->type, initExpr, ctx, sym->pos); - } -#endif else { // For non-static variables, allocate storage on the stack @@ -529,7 +402,6 @@ DeclStmt::EmitCode(FunctionEmitContext *ctx) const { sym->parentFunction = ctx->GetFunction(); InitSymbol(sym->storagePtr, sym->type, initExpr, ctx, sym->pos); } - } } }