Реализация критической секции в CUDA

Я изменил критическую секцию, потому что статья о переполнении стека, на которую я смотрел, не совсем подходила для моей ситуации.

Вот код

#include <stdio.h>
#include <stdlib.h>
#include <stdbool.h>

struct MatchedPixelsStruct {
    int* matched_pixels;
    int size;
};

#define MAX_ERROR_LENGTH 256

#ifdef __cplusplus
extern "C" {
#endif

struct MatchedPixelsStruct run();

#ifdef __cplusplus
}
#endif

__device__ int idx_iter = 0;

__global__ void processPixelsKernel( struct MatchedPixelsStruct *d_best_match_pixels, int *d_pixels ) {
    int len = 5;

    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    // каждый поток обрабатывает одну уникальную комбинацию left_right и down_up. left_right и down_up варьируются от -20 до 19.
    int left_right = (idx / 40) - 20;
    int down_up = (idx % 40) - 20;

    if ( idx > 1600 ) {
        return;
    }

    int* thread_matched_pixels2 = (int*) malloc( sizeof(int) * len );
    for (int i = 0; i < len; i++ ) {
        thread_matched_pixels2[i] = -1;
    }
    int cur_matched_score = 0;
    for ( int i = 0; i < len; i++ ) {
        int cur_score = left_right + down_up + d_pixels[i];
        if ( cur_score >= 0 ) {
            cur_matched_score += cur_score;
            thread_matched_pixels2[i] = cur_score;
        }
    }

    bool goon = true;
    while (goon) {

        if ( idx == idx_iter ) {
            // критическая секция
            printf("idx %d вошел\n", idx);
            if ( d_best_match_pixels->size < cur_matched_score ) {
                for ( int i = 0; i < len; i++ ) {
                    d_best_match_pixels->matched_pixels[i] = thread_matched_pixels2[i];
                }
                d_best_match_pixels->size = cur_matched_score;
            }
            goon = false;
            atomicAdd(&idx_iter, 1 );
        }
    }

}

struct MatchedPixelsStruct run( ) {
    int len = 5;
    int *pixels = (int*) malloc( sizeof(int) * len );
    for ( int i = 0; i < len; i++ ) {
        pixels[i] = i;
    }
    int *d_pixels;
    cudaMalloc(&d_pixels, sizeof(int) * len);
    cudaMemcpy(d_pixels, pixels, sizeof(int) * len, cudaMemcpyHostToDevice );

    struct MatchedPixelsStruct best_match_pixels;
    best_match_pixels.size = 0;

    best_match_pixels.matched_pixels = (int*)malloc(sizeof(int) * len);

    struct MatchedPixelsStruct *d_best_match_pixels;
    cudaMalloc(&d_best_match_pixels, sizeof(struct MatchedPixelsStruct));

    int* d_matched_pixels;
    cudaMalloc(&d_matched_pixels, sizeof(int) * len);
    cudaMemcpy(d_matched_pixels, best_match_pixels.matched_pixels, sizeof(int) * len, cudaMemcpyHostToDevice );
    cudaMemcpy(&(d_best_match_pixels->matched_pixels), &d_matched_pixels, sizeof(int*), cudaMemcpyHostToDevice);
    cudaMemcpy(&(d_best_match_pixels->size), &(best_match_pixels.size), sizeof(int), cudaMemcpyHostToDevice);

    // Параметры запуска ядра
    int blockSize = 256;
    int total_iterations = ( 20 * 2 ) * ( 20 * 2 ); // от 0 до 20 для перемещения пикселей. есть положительное и отрицательное движение. есть движение left_right и up_down.
    int numBlocks = (total_iterations + blockSize - 1) / blockSize;

    // Запуск ядра
    processPixelsKernel<<<numBlocks, blockSize>>>( d_best_match_pixels, d_pixels );

    cudaError_t err = cudaGetLastError();
    if (err != cudaSuccess) {
        printf("Ошибка CUDA: %s\n", cudaGetErrorString(err));
        exit(1);
    }

    cudaDeviceSynchronize();

    int size;
    cudaMemcpy(&size, &d_best_match_pixels->size, sizeof(int), cudaMemcpyDeviceToHost);
    int* matched_pixels = (int*) malloc( sizeof(int) * len );
    cudaMemcpy(matched_pixels, d_matched_pixels, sizeof(int) * len, cudaMemcpyDeviceToHost);

    best_match_pixels.size = size;
    best_match_pixels.matched_pixels = matched_pixels;

    cudaFree(d_best_match_pixels);
    cudaFree(d_matched_pixels);

    return best_match_pixels;
}

Часть критической секции гарантирует, что потоки входят в критическую секцию один за другим, начиная с потока 0 до потока 1600 последовательно. Но это работает только до 255, затем зависает.

Если я помещаю оператор printf внутрь цикла, чтобы увидеть, на каком значении idx_iter зависли ожидающие потоки, то зависание не происходит, но потоки с 0 по 30 отсутствуют, и от потока 31 происходит вход в критическую секцию, и это не зависает.

Вы не могли бы помочь мне понять, почему происходит такое странное поведение?