implementing CPU postprocessor

[MartinPulec]

refers to #75 (comment)

For the CPU postprocessor, you should already have data in decoder->coder->component[0-2].d_data. What the postprocessor actually does is in the conversion from the JPEG raw data to the output format. The data is almost always 8-bit YCbCr BT.601 and it may subsampled.

Let say that you want to decode to RGB. Your JPEG (in the original post) is subsampled 4:4:4, which is great, because you won't need to deal with the subsampling. So you'll just need to convert YCbCr->RGB, so I believe that in this easiest case it could be just something like (using coefficient from wiki), something similar to following:

const int height = decoder->coder->param_image.height;
const int width = decoder->coder->param_image.width;
for (int y = 0; y < height, y++) {
    for (int x = 0; x < width; x++) {
        int Y = coder->component[0].d_data[y * width + x];
        int U = coder->component[1].d_data[y * width + x] - 128;
        int V = coder->component[2].d_data[y * width + x] - 128;
        int R = Y + 1.402 * V;
        int G = Y - 0.344136 * U - 0.714136 * V;
        int B = Y 1.772 * U;
        coder->d_data_raw[3 * (y * width + x)] = R;
        coder->d_data_raw[3 * (y * width + x) + 1] = G;
        coder->d_data_raw[3 * (y * width + x) + 2] = B;
    }
}

There may be multiple aspects, anyways, except than subsampling. Eg. if the image dimension size not not divisible by 8. GPUJPEG also supports more output formats than RGB (it is specified by decoder->coder->param_image) etc.

UPDATE 2024-01-17: fixed d_data_raw which has been mistakenly d_data_row + pointers (either decoder->coder. or coder-> but not decoder->coder->)

[anthonyliot]

Thanks for the code let me try that in my fork to see if I am able to get it work on OSX :)

[anthonyliot]

[anthonyliot]

Adapting the code you provide I am able to decode with some artefact in the image but running on CPU

[MartinPulec]

I believe that the artifacts might have arisen in Huffman decoder. Actually, you can enforce Huffman decoding on CPU while leaving the rest on GPU just by enforcing this condition.

Just FYI, the dequantized data (before the postprocessor) could be possible to display with:
ffplay -f rawvideo -pixel_format yuv444p -video_size 1920x1080 output.yuv (for 444 divisible by 8 in both dimensions; if the planes are stored one right after another in the output file). But this is just for the record, it doesn't seem to be the case.

[anthonyliot]

Mmmh the code is using mainly C except for the cuda code :)
Typically the template is not supported by C, are you ok if we mixed C and C++ to enable the template support ?

[anthonyliot]

I made another push in order to implement the post processor and the preprocessor in CPU for now.
The code remain exactly similar to cuda with few small differences : 39d28a2

Overall I think it's closed but I have another issue now with the output, I must have a mistake in my logic over the kernel when converting to CPU:

[MartinPulec]

are you ok if we mixed C and C++ to enable the template support ?

If you mean the pre/postprocessor code reimplementation for CPU, I am totally fine with it, especially if it is in new files. We can discuss also for existing files; it has even been so between a68b3ec and 4d6aa8c (because MSVC used to poorly support C99).

[MartinPulec]

but I have another issue now with the output

Let me know if I can help you with something. As I am not the original author I may or may not know but at least some parts of code I've seen few times.

[anthonyliot]

Right now, I need to figure out why the code output for incorrect Raw data.

My suspicion is during the replacement of some cuda memcpy I make a mistake on the stride / width potentially that will explain this artefact.

If you look at my commit maybe you will see something.

[MartinPulec]

Hi, it seems that the values overflow/underflow sometimes, if you look at my snippet above (sorry that there were small mistakes, I've done it blindly), it behaves similarly as your new CPU postprocessor, however, if I wrap the computed RGB values with GPUJPEG_CLAMP, it seems to work:

int R = GPUJPEG_CLAMP(Y + 1.402 * V, 0, 255);
int G = GPUJPEG_CLAMP(Y - 0.344136 * U - 0.714136 * V, 0, 255);
int B = GPUJPEG_CLAMP(Y + 1.772 * U, 0, 255);

So in this light, I believe that there is a problem with int4 compat, which must have int32_t, not uint32_t in case it underflows.

[anthonyliot]

Hi @MartinPulec I just pushed some more change and I think the decoder is working now on CPU.
I was starting to look at the encoder, but there is one missing function gpujpeg_dct_cpu, any suggestion ??
Should I try to do exactly what I did for the preprocessor and postprocessor taking the cuda code and try to convert in c++ ?

The decoder give me a clean output now for the raw.rgb file:

[MartinPulec]

Should I try to do exactly what I did for the preprocessor and postprocessor taking the cuda code and try to convert in c++ ?

I think it is up to you how you decide. I believe that there are some ready-to-use CPU implementations as well, let say in libjpeg. Adapting the GPU code 1:1 may not be the most efficient solution from performance point of view but I think it is possible.

[anthonyliot]

Hi @MartinPulec I also prepared the code for runtime change of the encoder/decoder
Typically be able in your code to choose a gpujpeg on different gpujpeg_device
My idea will be to modify the function

GPUJPEG_API int
gpujpeg_init_device(int device_id, int flags);

to

GPUJPEG_API int
gpujpeg_init_device(gpujpeg_device device_id, int flags);

Where gpujpeg_device can be created like that

gpujpeg_device device("cuda:0")
gpujpeg_device device("opencl:0")
gpujpeg_device device("cpu")

6077209#diff-6b0299835ddf4a55e2f46a9de6c72b171531cfad7ffb9bf756f7f18705f06520

[MartinPulec]

Ok, why not. Anyways, I prefer not to introduce C++ to GPUJPEG API – the current code can be directly used from both C and C++, with C++ API it can be used only from C++ without a wrapper.

[anthonyliot]

Hi Martin how are you ?

So I am trying to really not use your CPU implementation but instead be able to generalize your kernel so I can use it on CPU.
The idea behind it's to have a real similar implementation almost identical across Acceleration CPU / CUDA / OPENCL

The other advantage will be for CPU to implement a similar kernel and later improve the speed by using openmp and multithread.
Unfortunately I am struggling with the implementation for the IDCT, here my current implementation copied from gpu and try to make it work on CPU.

In orde to be similar I right now pass the block / thread as parameter of the function.
All that needs to be cleanup but right now I just want to get a valid output.

Right now the raw data I get is full gray

Do you mind check the code to see what I missed please ? Right now in order to replace __syncthreads I am separating the kernel in three step (with openmp we will use an equivalent barrier)

Thanks

uint16_t gpujpeg_idct_cpu_quantization_table[64];

void
gpujpeg_idct_cpu_kernel_inplace(float* V8)
{
	//costants which are used more than once
	const float koeficient[6] = {0.4142135623f, 0.3535533905f, 0.4619397662f, 0.1989123673f, 0.7071067811f, -2.0f};
	
	V8[2] *= 0.5411961f;
	V8[4] *= 0.509795579f;
	V8[5] *= 0.601344887f;
	
	V8[1] = (V8[0] - V8[1]) * koeficient[1];
	V8[0] = V8[0] * koeficient[4] - V8[1];

	V8[3] = V8[2] * koeficient[1] + V8[3] * koeficient[2];
	V8[2] = V8[3] * koeficient[0] - V8[2];

	V8[6] = V8[5] * koeficient[2] + V8[6] * koeficient[0];
	V8[5] = -0.6681786379f * V8[6] + V8[5];

	V8[7] = V8[4] * koeficient[3] + V8[7] * 0.49039264f;
	V8[4] = V8[7] * koeficient[3] - V8[4];

	//instead of float tmp = V8[1]; V8[1] = V8[2] + V8[1]; V8[2] = tmp - V8[2];
	//we use this two operations (with a use of a multiply-add instruction)
	V8[1] = V8[2] + V8[1];
	V8[2] = koeficient[5] * V8[2] + V8[1];

	V8[4] = V8[5] + V8[4];
	V8[5] = 2.0f * V8[5] - V8[4];

	V8[7] = V8[6] + V8[7];
	V8[6] = koeficient[5] * V8[6] + V8[7];

	V8[0] = V8[3] + V8[0];
	V8[3] = koeficient[5] * V8[3] + V8[0];

	V8[5] = V8[6] * koeficient[0] + V8[5];
	V8[6] = V8[5] * -koeficient[4] + V8[6];
	V8[5] = V8[6] * koeficient[0] + V8[5];

	V8[3] = V8[3] + V8[4];
	V8[4] = koeficient[5] * V8[4] + V8[3];

	V8[2] = V8[2] + V8[5];
	V8[5] = koeficient[5] * V8[5] + V8[2];

	V8[1] = V8[6] + V8[1];
	V8[6] = koeficient[5] * V8[6] + V8[1];

	V8[0] = V8[0] + V8[7];
	V8[7] = koeficient[5] * V8[7] + V8[0];
}

#define MIMIC_BLOCK_THREAD_KERNEL_BEGIN \
    for (int blockIdx_x = 0; blockIdx_x < block_count_x; blockIdx_x++) { \
        for (int blockIdx_y = 0; blockIdx_y < block_count_y; blockIdx_y++) { \
            for (int threadIdx_x = 0; threadIdx_x < thread_per_block_x; threadIdx_x++) { \
                for (int threadIdx_y = 0; threadIdx_y < thread_per_block_y; threadIdx_y++) { \
                    for (int threadIdx_z = 0; threadIdx_z < thread_per_block_z; threadIdx_z++) { \

#define MIMIC_BLOCK_THREAD_KERNEL_END \
                    } \
                } \
            } \
        } \
	}


void gpujpeg_idct_cpu_kernel(int16_t* source, uint8_t* result, int output_stride, uint16_t* quantization_table, int block_count_x, int block_count_y, int thread_per_block_x, int thread_per_block_y, int thread_per_block_z) {

	float data[GPUJPEG_IDCT_BLOCK_Z][8][GPUJPEG_IDCT_BLOCK_Y][GPUJPEG_IDCT_BLOCK_X + 1];

	MIMIC_BLOCK_THREAD_KERNEL_BEGIN
      
	// **********************************************************************
	// **********************************************************************
	// **********************************************************************
	//here the grid is assumed to be only in x - it saves a few operations; if a larger
	//block count is used (e. g. GPUJPEG_IDCT_BLOCK_Z == 1), it would need to be adjusted,
	//the blockIdx_x not to exceed 65535. In the current state this function is good 
	//enough for a 67.1 MPix picture (8K is 33.1 MPix)

	//the first block of picture processed in this thread block
	unsigned int picBlockNumber = (blockIdx_x) * GPUJPEG_IDCT_BLOCK_Y * GPUJPEG_IDCT_BLOCK_X * GPUJPEG_IDCT_BLOCK_Z;

	//pointer to the begin of data for this thread block
	int16_t* sourcePtr = (int16_t*) (source) + picBlockNumber * 8;

	//variables to be used later more times (only one multiplication here)
	unsigned int z64 = threadIdx_z * 64;
	unsigned int x8 = threadIdx_x * 8;

	//data copying global -> shared, type casting int16_t -> float and dequantization.
	//16b reading gives only 50% efectivity but another ways are too complicated 
	//so this proves to be the fastest way
	for (int i = 0; i < 8; i++) {
		data[threadIdx_z][i][threadIdx_x][threadIdx_y] = sourcePtr[x8
				+ threadIdx_y + i * GPUJPEG_IDCT_BLOCK_X * GPUJPEG_IDCT_BLOCK_Y + z64 * 8]
				* quantization_table[threadIdx_x * 8 + threadIdx_y];
	}

	MIMIC_BLOCK_THREAD_KERNEL_END
	
	// __syncthreads();

	float x[8];
	
	MIMIC_BLOCK_THREAD_KERNEL_BEGIN
	//kompilator delal hrozne psi kusy - zbytecne kopirovani konstant do
	//registru atp., bylo jednodussi napsat to v assembleru nez snazit se ho
	//presvedcit, aby nedelal blbosti; vsechny konstanty se pouzivaji primo
	//hodnotou, nestrkaji se zbytecne do registru

	//here the data are being processed by columns - each thread processes one column
	x[0] = data[threadIdx_z][threadIdx_x][0][threadIdx_y];
	x[1] = data[threadIdx_z][threadIdx_x][4][threadIdx_y];
	x[2] = data[threadIdx_z][threadIdx_x][6][threadIdx_y];
	x[3] = data[threadIdx_z][threadIdx_x][2][threadIdx_y];
	x[4] = data[threadIdx_z][threadIdx_x][7][threadIdx_y];
	x[5] = data[threadIdx_z][threadIdx_x][5][threadIdx_y];
	x[6] = data[threadIdx_z][threadIdx_x][3][threadIdx_y];
	x[7] = data[threadIdx_z][threadIdx_x][1][threadIdx_y];
	
	gpujpeg_idct_cpu_kernel_inplace(x);

	data[threadIdx_z][threadIdx_x][0][threadIdx_y] = x[0];
	data[threadIdx_z][threadIdx_x][1][threadIdx_y] = x[1];
	data[threadIdx_z][threadIdx_x][2][threadIdx_y] = x[2];
	data[threadIdx_z][threadIdx_x][3][threadIdx_y] = x[3];
	data[threadIdx_z][threadIdx_x][4][threadIdx_y] = x[4];
	data[threadIdx_z][threadIdx_x][5][threadIdx_y] = x[5];
	data[threadIdx_z][threadIdx_x][6][threadIdx_y] = x[6];
	data[threadIdx_z][threadIdx_x][7][threadIdx_y] = x[7];

	MIMIC_BLOCK_THREAD_KERNEL_END

	MIMIC_BLOCK_THREAD_KERNEL_BEGIN

	x[0] = data[threadIdx_z][threadIdx_x][threadIdx_y][0];
	x[1] = data[threadIdx_z][threadIdx_x][threadIdx_y][4];
	x[2] = data[threadIdx_z][threadIdx_x][threadIdx_y][6];
	x[3] = data[threadIdx_z][threadIdx_x][threadIdx_y][2];
	x[4] = data[threadIdx_z][threadIdx_x][threadIdx_y][7];
	x[5] = data[threadIdx_z][threadIdx_x][threadIdx_y][5];
	x[6] = data[threadIdx_z][threadIdx_x][threadIdx_y][3];
	x[7] = data[threadIdx_z][threadIdx_x][threadIdx_y][1];

	gpujpeg_idct_cpu_kernel_inplace(x);

	//output will be written by 8B (one row) which is the most effective way
	uint64_t tempResult;
	uint64_t* tempResultP = &tempResult;

	for (int i = 0; i < 8; i++) {
		//this would be faster but will work only for 100% quality otherwise some values overflow 255
		//((uint8_t*) tempResultP)[i] = __float2uint_rz(x[i] + ((float) 128.0));

		//cast float to uint8_t with saturation (.sat) which cuts values higher than 
		//255 to 255 and smaller than 0 to 0; cuda can't use a reg smaller than 32b 
		//(though it can convert to 8b for the saturation purposes and save to 32b reg)
		int save = rintf(x[i] + 128.0F);
		save = save < 0 ? 0 : save > 255 ? 255 : save;
		((uint8_t*) tempResultP)[i] = save;
	}

	//between data writing and sync it's good to compute something useful 
	// - the sync will be shorter.
	unsigned int picBlockNumber = (blockIdx_x) * GPUJPEG_IDCT_BLOCK_Y * GPUJPEG_IDCT_BLOCK_X
			* GPUJPEG_IDCT_BLOCK_Z;

	//variables to be used later more times (only one multiplication here)
	unsigned int z64 = threadIdx_z * 64;
	unsigned int x8 = threadIdx_x * 8;

	//output pointer (the begin for this thread block)
	unsigned int firstByteOfActualBlock = x8 + z64 + picBlockNumber;

	uint8_t* resultPtr = ((uint8_t*) result) + firstByteOfActualBlock
			+ (threadIdx_y + ((firstByteOfActualBlock / output_stride) * 7))
					* output_stride;

	//writing result - one row of a picture block by a thread
	*((uint64_t*) resultPtr) = tempResult;

	MIMIC_BLOCK_THREAD_KERNEL_END
}

int gpujpeg_idct_cpu(struct gpujpeg_decoder* decoder)
{
    struct gpujpeg_coder* coder = &decoder->coder;

    // Encode each component
    for ( int comp = 0; comp < coder->param_image.comp_count; comp++ ) {
        // Get component
        struct gpujpeg_component* component = &coder->component[comp];

        int roi_width = component->data_width;
        int roi_height = component->data_height;
        assert(GPUJPEG_BLOCK_SIZE == 8);

        int block_count_x = roi_width / GPUJPEG_BLOCK_SIZE;
        int block_count_y = roi_height / GPUJPEG_BLOCK_SIZE;

        // Get quantization table
        uint16_t* d_quantization_table = decoder->table_quantization[decoder->comp_table_quantization_map[comp]].d_table;
		
		memcpy((void*)gpujpeg_idct_cpu_quantization_table, d_quantization_table, 64 * sizeof(uint16_t));

        gpujpeg_idct_cpu_kernel(
            component->d_data_quantized,
            component->d_data,
            component->data_width,
            d_quantization_table,
			gpujpeg_div_and_round_up(block_count_x * block_count_y,
				(GPUJPEG_IDCT_BLOCK_X * GPUJPEG_IDCT_BLOCK_Y * GPUJPEG_IDCT_BLOCK_Z) / GPUJPEG_BLOCK_SIZE), 1,
				GPUJPEG_IDCT_BLOCK_X, GPUJPEG_IDCT_BLOCK_Y, GPUJPEG_IDCT_BLOCK_Z
        );
    }

    return 0;
}

Here the output with this code - full grey :(

[MartinPulec]

Hi @anthonyliot, sorry for the late reply, I've currently too much tasks under my hands.

I am afraid that I'm going to disappoint you, but I am not able to help you right now. The thing is that I am not very familiar with that code and I actually consider it over-optimized (although it is fast, so I've merged it; but it doesn't need to be so for the CPU version). It will take me too much time and I don't think it is worth torture oneself if the current CPU code is working (?).

[anthonyliot]

Hi @MartinPulec no worry, I am also will a lot of task right now and was not able to make some progress.
I think it is important in a way to understand for me the behavior and why it is not working on CPU to help on the transformation to OpenCL.

But I start understanding more of what I was missing. My last update on the change I did I try to use a more C approach syntax and modify slightly the code so instead to pass a device to the gpujpeg_init as int for the cuda gpu number, now you can specify the gpujpeg_device who can be
cuda:0 cuda:1 ... cl:0 cl:1 ... cpu

But lot of work remain to do :) I will try to find some few hours this week to make some progress

[MartinPulec]

I think it is important in a way to understand for me the behavior and why it is not working on CPU to help on the transformation to OpenCL.

I totally understand.

Not sure if it helps you – just from curiosity I've just tried to compare outcome of your code with the original, and for very simple example (just 2 DCT values in a single 8x8 block) and the values look the same. You can try out tst.cu.txt – it is basically your code with gpujpeg_idct_cpu transformed to main function without GPUJPEG dependencies + imported the appropriate GPU kernel from gpujpeg_dct_gpu.cu.

[anthonyliot]

Hi @MartinPulec I am working on the codebase, I create a PR on my side so you can easily the change I am doing and get your feedback : anthonyliot#1

The code still need a lot of cleanup and improvement but the idea is mainly:

You need to have a gpujpeg_device as parameter to initialize gpujpeg the device can be 'cpu', 'opencl:0', 'cuda:0' ....
Right now the code is only building on OSX for CPU only, I just wanted to start reorganizing.

I have one class accelerate with a list of pointer function per acceleration, right now the CPU decoder is implemented (But output is incorrect) and will start adding back cuda decoder.

If that work then the logic for the encoder should be the same.
I am also using an approach to use the exact same codebase kernel for both CPU and CUDA (need to check how that will work with opencl)

The interesting part is potentially when everything work the CPU could be just modify a little bit so I can use OpenMP and then all the for loop in the kernel to mimic cuda will be accelerated also on CPU using the multicore.

There is still a lot to do and right now the output of my CPU decode still give me weird output :

In case you have an idea why it look like that do not hesitate to share :)

[anthonyliot]

The timer is also implemented to support different acceleration (The original version was using cudaEvent):

GPUJPEG rev a45db89

Decoding Image [/Volumes/SABRENT/github/imageloader/images/test.jpg] to 444-u8-p012 RGB
Load Image:                0.04 ms
[GPUJPEG] [Warning] JPEG data contains not supported APP1 marker
[GPUJPEG] [Debug] coder image reconfiguration

Allocation Info:
    Segment Count:            1
    Allocated Data Size:      304x232
    Raw Buffer Size:          0.2 MiB
    Preprocessor Buffer Size: 0.2 MiB
    DCT Buffer Size:          0.4 MiB
    Compressed Buffer Size:   1.6 MiB
    Huffman Temp buffer Size: 1.6 MiB
    Structures Size:          0.4 KiB
    Total GPU Memory Size:    4.0 MiB

Using faster kernel for postprocessor (precompiled 1x1, 1x1, 1x1).
 -Stream Reader:         0.1570 ms
 -Copy To Device:        0.0640 ms
 -Huffman Decoder:       1.5960 ms
 -DCT & Quantization:    3.4220 ms
 -Postprocessing:        2.8250 ms
 -Copy From Device:      0.0160 ms
Decode Image Device:     6.2470 ms (only in-device 'cpu' processing)
Decode Image Bare:       8.0041 ms (without copy to/from 'cpu' memory)
Decode Image:            8.0841 ms
[GPUJPEG] [Warning] JPEG data contains not supported APP1 marker
Save Image:              0.38 ms
Decompressed Size:       202500 bytes [/Volumes/SABRENT/github/imageloader/images/test.raw]
Program ended with exit code: 0

[anthonyliot]

Hi @MartinPulec

I did some progress I was able to get the code (Only decoder) fully working with cuda on my linux machine and output correct result. All using this new acceleration system.

I know something is wrong in my IDCT kernel, by modifying how I try to mimic the running kernel in CPU I am able to get an output but with a lot of noise / artefact

Here for example the result of the command -L on the gpupeg (list of devices) need to complete some value for CPU and OpenCL

$ ./gpujpegtool -L
GPUJPEG rev b7fdbec

There is 1 device supporting CPU
Device #0: "cpu"
  Compute capability: 0.0
  Total amount of global memory: 0 KiB
  Total amount of constant memory: 0 KiB
  Total amount of shared memory per block: 0 KiB
  Total number of registers available per block: 0
  Multiprocessors: 0

There is 1 device supporting CUDA
Device #0: "NVIDIA GeForce GTX 1070"
  Compute capability: 0.0
  Total amount of global memory: 8278592 KiB
  Total amount of constant memory: 64 KiB
  Total amount of shared memory per block: 48 KiB
  Total number of registers available per block: 65536
  Multiprocessors: 15

There is 1 device supporting OPENCL
Device #0: "NVIDIA GeForce GTX 1070"
  Compute capability: 0.0
  Total amount of global memory: 8278592 KiB
  Total amount of constant memory: 64 KiB
  Total amount of shared memory per block: 48 KiB
  Total number of registers available per block: 0
  Multiprocessors: 15

To decode an image you have to just do :

./gpujpegtool -v -d ../../imageloader/images/test.jpg test_cuda.raw -D cuda
./gpujpegtool -v -d ../../imageloader/images/test.jpg test_cuda.raw -D cuda:0
./gpujpegtool -v -d ../../imageloader/images/test.jpg test_cuda.raw -D cpu

[anthonyliot]

Ok I got the decode working on CPU the issue is related to a misunderstanding on the __syncthread behavior.
after fixing the way I use it in the kernel we have the decode CPU and GPU working now.
The decode is missing some implementation but I will work on the encoder now

[anthonyliot]

Hi @MartinPulec
I added the Encoding part but it seems I must have made a mistake because the output (using cuda) is not giving me a valid output jpeg.
The decoder CPU and GPU work fine, but I try to encode the RAW data to JPEG the result is using cuda is what I obtains now.
Also just wanted to be sure I see two implementation of the encoder one for cuda == 1 and cuda == 2, it seems the code for the cuda == 1 is broken Am I correct ?/home/anthonyliot/Desktop/github/gpujpeg/build/test_cuda.jpg

Anyway I attached the output result from CUDA encoder in case you have an idea what can give me such result

[anthonyliot]

Hi @MartinPulec,

I cleanup the branch I am working it and rebase on top of the latest change in your main branch.

Right now I am able to get the full Cuda pipeline working back Encoder / Decoder with and without restart interval.
For the CPU I have the decoder but the encoder give me some trouble with and without restart interval, it means I have some trouble with some kernel implementation (preprocessor and dct potentially) because using restart 0 I am keep using the cpu encoder from gpujpeg.

Here the output with restart interval 0 using CPU - I have more debugging work to do :)

[anthonyliot]

HI @MartinPulec
I was able to make some progress on my debug of the CPU encoder using the same approach as cuda.
The kernel use return on the DCT and on CPU I cannot interrupt the loop using return its more like a break in this case.

After fixing that I was ablate get a better output (Still using the interval 0 - like that I can focus first on the DCT and Preprocessor kernel)

The output is almost good (we see the image) but there is still something wrong, any idea ?

[anthonyliot]

I see in the code something related to alignment :

union PackedInteger
{
    struct __align__(8)
    {
        int16_t hShort1 /*__attribute__((aligned(8)));*/
        int16_t hShort2 /*__attribute__((aligned(8)));*/
    };
    int32_t hInt;
};

I do not see that being used could it be something related to that I am missing ?

[anthonyliot]

Found the issue it was related to some syncthread on the shared array, but that say I still feel there is an issue in my DCT kernel because the output JPEG seems having some weird vertical line.

I need to fix the DCT kernel and the ENCODER kernel then I will be able to start OpenCL,
Maybe I will initiate some work next week to do the decoder using OpenCL