square Матрица транспонирована с CUDA

Question

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

Почему мой массив не переносится правильно? Кто-нибудь может мне помочь или дать мне какой-нибудь намек? Я буду признателен. Большое спасибо!

есть весь код процессора:

__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 <=w && yIdx <=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(&d_idata, mem_size)); 
    checkCuda(cudaMalloc(&d_cdata, mem_size)); 
    // host 
    for (int j = 0; j < nx; j++) 
     for (int i = 0; i < 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(&startEvent)); 
    checkCuda(cudaEventCreate(&stopEvent)); 
    float ms; 
    checkCuda(cudaEventRecord(startEvent, 0)); 
    transpose_naive<<<dimGrid, dimBlock>>>(d_cdata, d_idata,nx,nx); 
    checkCuda(cudaEventRecord(stopEvent, 0)); 
    checkCuda(cudaEventSynchronize(stopEvent)); 
    checkCuda(cudaEventElapsedTime(&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"); 
} 

Answer 1

Единственный недостаток в коде неправильные связанные проверки в следующей строке ядра.

if (xIdx <=w && yIdx <=h) { 

Как индексы от 0 к w-1 и 0 к h-1 для й и у размеры соответственно, if условия должно быть следующим:

if (xIdx <w && yIdx <h) { 

Answer 2

Я думаю, что есть что-то не так с выходом файла «i.txt» и «t.txt», иначе программа выглядит правильно. Я внес некоторые незначительные изменения в свой код, добавив проверку и печать ошибок в стандартный поток вывода. Я печатаю последнюю (1020 - 1024) матрицу 3 x 3, чтобы перекрестно проверить транспонирование. Запустите его в своей системе и проверьте правильность преобразования матрицы?

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

#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); } 
inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort = true) 
{ 
if (code != cudaSuccess) 
{ 
    fprintf(stderr, "GPUassert: %s %s %d\n", cudaGetErrorString(code),file, line); 
    if (abort) exit(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 <= w && yIdx <= 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; 
gpuErrchk(cudaMalloc(&d_idata, mem_size)); 
gpuErrchk(cudaMalloc(&d_cdata, mem_size)); 
// host 
for (int j = 0; j < nx; j++) 
    for (int i = 0; i < nx; i++) 
     h_idata[j*nx + i] = j*nx + i; 

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

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

printf("the time %5f ", ms); 
printf("\n"); 
for (int i = 1020; i < 1024; i++) { 
    for (int j = 1020; j < 1024; j++) { 
    printf("%.2f ", h_idata[i*nx + j]); 
     } 
     printf("\n"); 
} 

printf("\n"); 
for (int i = 1020; i < 1024; i++) { 
    for (int j = 1020; j < 1024; j++) { 
     printf("%.2f ", h_cdata[i*nx + j]); 
    } 
    printf("\n"); 
} 
//savetofile(h_idata, "i.txt", nx, nx); 
//savetofile(h_cdata, "t.txt", nx, nx); 
//error_exit: 
// cleanup 
gpuErrchk(cudaEventDestroy(startEvent)); 
gpuErrchk(cudaEventDestroy(stopEvent)); 
gpuErrchk(cudaFree(d_cdata)); 
gpuErrchk(cudaFree(d_idata)); 
free(h_idata); 
free(h_cdata); 
//system("pause"); 
}