转置的好用的cuda程序

通过sample的例子自己改编的一个例子

#include <stdio.h>
#define BLOCK_DIM 5

// Transpose kernel (see transpose CUDA Sample for details)
__global__ void d_transpose(float *odata, float *idata, int width, int height)
{
    __shared__ float block[BLOCK_DIM][BLOCK_DIM+1];

    // read the matrix tile into shared memory
    unsigned int xIndex = blockIdx.x * BLOCK_DIM + threadIdx.x;
    unsigned int yIndex = blockIdx.y * BLOCK_DIM + threadIdx.y;

    if ((xIndex < width) && (yIndex < height))
    {
        unsigned int index_in = yIndex * width + xIndex;
        block[threadIdx.y][threadIdx.x] = idata[index_in];
    }

    __syncthreads();

    // write the transposed matrix tile to global memory
    xIndex = blockIdx.y * BLOCK_DIM + threadIdx.x;
    yIndex = blockIdx.x * BLOCK_DIM + threadIdx.y;

    if ((xIndex < height) && (yIndex < width))
    {
        unsigned int index_out = yIndex * height + xIndex;
        odata[index_out] = block[threadIdx.x][threadIdx.y];
    }
}
void print_arr(float a[],int row,int col,char * info){
        printf("%s\n",info);
        for(int i=0;i<row;i++){
                for(int j=0;j<col;j++){
                        printf("%f ",a[i*col+j]);
                }
                printf("\n");
        }
}
int iDivUp(int a, int b)
{
    return (a % b != 0) ? (a / b + 1) : (a / b);
}

/*
    Transpose a 2D array (see SDK transpose example)
*/
extern "C"
void transpose(float *d_src, float *d_dest,int width, int height)
{
    dim3 grid(iDivUp(width, BLOCK_DIM), iDivUp(height, BLOCK_DIM), 1);
    dim3 threads(BLOCK_DIM, BLOCK_DIM, 1);
    d_transpose<<< grid, threads >>>(d_dest, d_src, width, height);
}

int main(){
	const int  nx = 32;
  	const int  ny = 32;
  	const int mem_size = nx*ny*sizeof(float);

	float *h_idata = (float *)malloc(mem_size);
  	float  *h_cdata = (float *)malloc(mem_size);
  	float *h_tdata = (float*)malloc(mem_size);

	float *d_idata, *d_cdata, *d_tdata;
  	cudaMalloc(&d_idata, mem_size) ;
 	cudaMalloc(&d_cdata, mem_size) ;
  	cudaMalloc(&d_tdata, mem_size) ;

	for (int j = 0; j < ny; j++){
    		
			for (int i = 0; i < nx; i++){
				h_idata[j*nx + i] = j+0.1;//j*nx + i;
			}
		}	
	
	cudaMemcpy(d_idata, h_idata, mem_size, cudaMemcpyHostToDevice) ;;

	
	transpose(d_idata,d_tdata ,nx,ny);

	cudaMemcpy(h_tdata, d_tdata, mem_size, cudaMemcpyDeviceToHost);
	
	print_arr(h_idata,nx,ny,"origin data is");	
	print_arr(h_tdata,nx,ny,"transposed data is");	

	cudaFree(d_tdata) ;
  	cudaFree(d_cdata) ;
  	cudaFree(d_idata) ;
	free(h_idata);
  	free(h_tdata);
  	free(h_cdata);
	return 0;
}

还有一个程序，挺奇怪的现在还不知道怎么用，有知道的可以交流一下

#include <stdio.h>
#include <cuda.h>
const int TILE_DIM = 16;
const int BLOCK_ROWS = 8;
void print_arr(float a[],int row,int col,char * info){
        printf("%s\n",info);
        for(int i=0;i<row;i++){
                for(int j=0;j<col;j++){
                        printf("%f ",a[i*col+j]);
                }
                printf("\n");
        }
}
__global__ void transposeNaive(float *odata, const float *idata)
{
  int x = blockIdx.x * TILE_DIM + threadIdx.x;
  int y = blockIdx.y * TILE_DIM + threadIdx.y;
  int width = gridDim.x * TILE_DIM;

  for (int j = 0; j < TILE_DIM; j+= BLOCK_ROWS)
    odata[x*width + (y+j)] = idata[(y+j)*width + x];
}
int main(){
	const int nx = 32;
  	const int ny = 32;
  	const int mem_size = nx*ny*sizeof(float);

	float *h_idata = (float*)malloc(mem_size);
  	float *h_cdata = (float*)malloc(mem_size);
  	float *h_tdata = (float*)malloc(mem_size);

	float *d_idata, *d_cdata, *d_tdata;
  	cudaMalloc(&d_idata, mem_size) ;
 	cudaMalloc(&d_cdata, mem_size) ;
  	cudaMalloc(&d_tdata, mem_size) ;

	for (int j = 0; j < ny; j++)
    		for (int i = 0; i < nx; i++)
      			h_idata[j*nx + i] = j%32;//j*nx + i;	
	
	cudaMemcpy(d_idata, h_idata, mem_size, cudaMemcpyHostToDevice) ;

	dim3 dimGrid(nx/TILE_DIM, ny/TILE_DIM, 1);
  	dim3 dimBlock(TILE_DIM, BLOCK_ROWS, 1);
	cudaMemset(d_tdata, 0, mem_size) ;
	transposeNaive<<<dimGrid, dimBlock>>>(d_tdata, d_idata);

	cudaMemcpy(h_tdata, d_tdata, mem_size, cudaMemcpyDeviceToHost);
	
	print_arr(h_idata,32,32,"origin data is");	
	print_arr(h_tdata,32,32,"transposed data is");	

	cudaFree(d_tdata) ;
  	cudaFree(d_cdata) ;
  	cudaFree(d_idata) ;
	free(h_idata);
  	free(h_tdata);
  	free(h_cdata);
	return 0;
}