square Matrix transpose with CUDA
<p>I'm trying to write  the  matrix transpose algorithm. I test this program with matrix size equal to 1024, the result shows that not all elements are in the right places.</p>

<p>Why isn't my array transposing correctly? Does anyone can help me or give me any hint? I will appreciate it. Thanks a lot!</p>

<p>there is the whole cpu code:</p>

<pre><code>__global__ void transpose_naive (float *out, float *in, int w, int h )
{
    unsigned int xIdx = blockDim.x * blockIdx.x + threadIdx.x;
    unsigned int yIdx = blockDim.y * blockIdx.y + threadIdx.y;
    if ( xIdx &lt;=w &amp;&amp; yIdx &lt;=h ) {
        unsigned int idx_in = xIdx + w * yIdx;
        unsigned int idx_out = yIdx + h * xIdx;
        out[idx_out] = in[idx_in];
    }
}

int main()

{
    int nx=1024;
    int mem_size = nx*nx*sizeof(float);
    int t=32;
    dim3 dimGrid(((nx-1)/t) +1, ((nx-1)/t) +1);
    dim3 dimBlock(t,t);

    float *h_idata = (float*)malloc(mem_size);
    float *h_cdata = (float*)malloc(mem_size);
    float *d_idata, *d_cdata;
    checkCuda(cudaMalloc(&amp;d_idata, mem_size) );
    checkCuda(cudaMalloc(&amp;d_cdata, mem_size) );   
    // host
    for (int j = 0; j &lt; nx; j++)
        for (int i = 0; i &lt; nx; i++)
            h_idata[j*nx + i] = j*nx + i;

    // device
    checkCuda(cudaMemcpy(d_idata, h_idata, mem_size, cudaMemcpyHostToDevice) );

    // events for timing
    cudaEvent_t startEvent, stopEvent;
    checkCuda(cudaEventCreate(&amp;startEvent) );
    checkCuda(cudaEventCreate(&amp;stopEvent) );
    float ms;
    checkCuda( cudaEventRecord(startEvent, 0) );
    transpose_naive&lt;&lt;&lt;dimGrid, dimBlock&gt;&gt;&gt;(d_cdata, d_idata,nx,nx);
    checkCuda(cudaEventRecord(stopEvent, 0) );
    checkCuda(cudaEventSynchronize(stopEvent) );
    checkCuda(cudaEventElapsedTime(&amp;ms, startEvent, stopEvent) );
    checkCuda( cudaMemcpy(h_cdata, d_cdata, mem_size, cudaMemcpyDeviceToHost) );

    printf("the time %5f ", ms);
    printf("\n");
    savetofile(h_idata,"i.txt",nx,nx);
    savetofile(h_cdata,"t.txt",nx,nx);

error_exit:
    // cleanup
    checkCuda(cudaEventDestroy(startEvent) );
    checkCuda(cudaEventDestroy(stopEvent) );
    checkCuda( cudaFree(d_cdata) );
    checkCuda( cudaFree(d_idata) );
    free(h_idata);
    free(h_cdata);
    system("pause"); 
}
</code></pre>