Алгоритм сортировки слиянием на CUDA

const int BLOCK_SIZE1 = 1024;
 
 
__global__ void merge_sort(double *DA)
{
    __shared__ double DAS1[BLOCK_SIZE1];
    __shared__ double DAS2[BLOCK_SIZE1];
    int tx = threadIdx.x;
    int i, j, k, m;
 
    DAS1[tx] = DA[tx];
    [COLOR="red"]__syncthreads();[/COLOR]
 
    for (m = 2; m <= BLOCK_SIZE1; m *= 2)
    {
        if (tx % m == 0)
        {
            i = tx; j = tx + m / 2; k = tx;
            while (i < tx + m / 2 || j < tx + m)
            {
                if (i < tx + m / 2 && j == tx + m)
                {
                    DAS2[k] = DAS1[i]; i++;  k++;
                }
                if (i == tx + m / 2 && j < tx + m)
                {
                    DAS2[k] = DAS1[j]; j++;  k++;
                }
                if (i < tx + m / 2 && j < tx + m)
                {
                    if (DAS1[i] <= DAS1[j])
                    {
                        DAS2[k] = DAS1[i]; i++;  k++;
                    }
                    else
                    {
                        DAS2[k] = DAS1[j]; j++;  k++;
                    }
                }
            }
        }
        __syncthreads();
 
        DAS1[tx] = DAS2[tx];
        __syncthreads();
    }
 
    DA[tx] = DAS1[tx];
    __syncthreads();
}
 
 
 
 
int main()
{
    double *DA;
    
 
    cudaMalloc((void**)&DA, BLOCK_SIZE1 * sizeof(double));
    thrust::host_vector<double> HA(BLOCK_SIZE1);
    thrust::generate(HA.begin(), HA.end(), rand() % 3);
    cudaMemcpy(DA, HA.data(), BLOCK_SIZE1 * sizeof(double), cudaMemcpyHostToDevice);
    [COLOR="Red"]merge_sort << < 1, BLOCK_SIZE1 >> > (DA);
[/COLOR]
    cudaFree(DA);
}

int tx = threadIdx.x;

int tx = blockDim.x * blockIdx.x + threadIdx.x;

for (m = 2; m <= BLOCK_SIZE1; m *= 2)

for (m = 2; m <= blockDim.x * gridDim.x; m *= 2)

__shared__ double DAS1[BLOCK_SIZE1];
__shared__ double DAS2[BLOCK_SIZE1];

__shared__ double DAS1[blockDim.x * gridDim.x];
__shared__ double DAS2[blockDim.x * gridDim.x];