CUDA实例练习（四）：矩阵转置

#include <stdio.h>
#include "gputimer.h"
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdlib.h>

const int N = 1024;
const int K = 32;

void fill_matrix(float * mat){
    for (int i = 0; i < N*N; i++)
        mat[i] = (float)i;
}

void print_matrix(float *mat)
{
    for (int j = 0; j < N; j++)
    {
        for (int i = 0; i < N; i++) { printf("%4.4g ", mat[i + j*N]); }
        printf("
");
    }
}

__global__ void transpose_serial(float in[], float out[]){
    for (int i = 0; i < N; i++)
        for (int j = 0; j < N; j++)
            out[i + j*N] = in[j + i*N];
}

__global__ void transpose_parallel_per_row(float in[], float out[]){
    int i = threadIdx.x;

    for (int j = 0; j < N; j++)
        out[j + i*N] = in[i + j*N];
}

__global__ void transpose_parallel_per_element(float in[], float out[]){
    int i = blockIdx.x * K + threadIdx.x;
    int j = blockIdx.y * K + threadIdx.y;
    out[j + i*N] = in[i + j*N];
}
int main(void){
    int numbytes = N * N * sizeof(float);

    float *in = (float *)malloc(numbytes);
    float *out = (float *)malloc(numbytes);
    fill_matrix(in);
    
    float *d_in, *d_out;

    cudaMalloc((void **)&d_in, numbytes);
    cudaMalloc((void **)&d_out, numbytes);
    cudaMemcpy(d_in, in, numbytes, cudaMemcpyHostToDevice);

    GpuTimer timer;
    timer.Start();
    transpose_serial << <1, 1 >> >(d_in, d_out);
    timer.Stop();
    cudaMemcpy(out, d_out, numbytes, cudaMemcpyDeviceToHost);
    printf("transpose_serial:%g ms.
", timer.Elapsed());

    timer.Start();
    transpose_parallel_per_row << <1, N >> >(d_in, d_out);
    timer.Stop();
    cudaMemcpy(out, d_out, numbytes, cudaMemcpyDeviceToHost);
    printf("transpose_parallel_per_row:%g ms.
", timer.Elapsed());

    dim3 blocks(N / K, N / K);
    dim3 threads(K, K);
    timer.Start();
    transpose_parallel_per_element << <blocks, threads >> >(d_in, d_out);
    timer.Stop();
    cudaMemcpy(out, d_out, numbytes, cudaMemcpyDeviceToHost);
    printf("transpose_parallel_per_element: %g ms.
", timer.Elapsed());

    cudaFree(d_in);
    cudaFree(d_out);

    return 0;
}