Forcing loop iterations to execute sequentially

You can force serial execution with loop unrolling. I don't know what your N and M values are, but have you tried doing partial or full unrolling?

Unless your kernel is being falsely compiled as NDRange, loops are never executed out of order in single work-item kernels. NDRange kernels have a scheduler that can change the order of threads at runtime, but single work-item kernels do not have n scheduler and loop iterations are guaranteed to be executed in-order. What is happening here is that your printfs are coming out out-of-order, not the iterations themselves. I am still not sure how printfs are implemented on the hardware, but there is likely some on-chip buffering in place and hence, I don't think correct ordering of printfs are guaranteed on the hardware. You should probably not use printf for debugging ordering on the hardware.

#define M 100
# define N 1000
//__attribute__((max_global_work_dim(0)))
__kernel void ast(__global float* restrict in, __global float* restrict out)
{
	__local float lmem;
	
	for (int i = 0; i < M; i++)
	{
		lmem = in;
	}
	for (uint outer = 0; outer < N; ++outer)
	{
		uint wr_bank_sel = outer & 0x1;
		uint rd_bank_sel = !(outer & 0x1);
		for (uint inner = 0; inner < M; ++inner)
		{
			lmem = lmem * 10.0f/(inner + outer); // placeholder math op but real dependencies
		}
	}
	
	for (int i = 0; i < M; i++)
	{
		out = lmem;
	}
}

The kernel is compiled for single work-item execution.
The kernel has a required work-group size of (1, 1, 1).
Loop Report:
 + Loop "Block1" (file test.cl line 9)
   Pipelined well. Successive iterations are launched every cycle.
   
   
 + Loop "Block2" (file test.cl line 14)
 | Pipelined with successive iterations launched every cycle.
 |  
 | Iterations executed serially across the region listed below.
 | Only a single loop iteration will execute inside the listed region.
 | This will cause performance degradation unless the region is pipelined well
 | (can process an iteration every cycle).
 |      
 |     Loop "Block3" (file test.cl line 18)
 |     due to:
 |     Memory dependency on Load Operation from: (file test.cl line 20)
 |       Store Operation (file test.cl line 20)
 | 
 | 
 |-+ Loop "Block3" (file test.cl line 18)
     Pipelined well. Successive iterations are launched every cycle.
     
     
 + Loop "Block5" (file test.cl line 24)
   Pipelined well. Successive iterations are launched every cycle.

Thanks HRZ, unfortunately I'm still having this issue. The test code you posted detects the serial execution dependency fine however I get very inconsistent results in the compiler for my program. I'm going to keep trying to create simplified versions to show you what I mean but if you have any# pragma recommendations or anything in the meantime let me know.

I am afraid I do not know of any pragmas that could force serial execution. Altera has many undocumented pargmas, they could also have one for this, but rarely anybody outside of Altera knows about them.