Не удается выполнить ядро ​​устройства в CUDA

Question

Я пытаюсь вызвать ядро ​​устройства в глобальном ядре. Мое глобальное ядро ​​- это умножение матрицы, а ядро ​​устройства находит максимальное значение и индекс в каждом столбце матрицы продукта. Ниже приводится код:

__device__ void MaxFunction(float* Pd, float* max) 
{ 
    int x = (threadIdx.x + blockIdx.x * blockDim.x); 
    int y = (threadIdx.y + blockIdx.y * blockDim.y); 
    int k = 0; 
    int temp = 0; int temp_idx = 0; 
    for (k = 0; k < wB; ++k) { 
    if(Pd[x*wB + y] > temp){ 
    temp = Pd[x*wB + y]; 
    temp_idx = x*wB + y; 
    } 
     max[y*2 + 0] = temp; 
     max[y*2 + 1] = temp_idx; 
    } 
} 

__global__ void MatrixMulKernel(float* Md, float* Nd, float* Pd, float* max) 
{ 
    // declare cache in the shared memory 
    __shared__ float Mds[blockD][blockD]; 
    __shared__ float Nds[blockD][blockD]; 

    float Pvalue = 0; 
    // Loop over the Md and Nd block dimension required to compute the Pd element 
    for (int m = (wA * blockD * blockIdx.y), n = (blockD * blockIdx.x); 
          m < ((wA * blockD * blockIdx.y)+wA-1); 
             m += blockD, n += (blockD*hB)){ 

    // collaboratively loading of Md and Nd blocks into shared memory  
    Mds[threadIdx.y][threadIdx.x] = Md[m + wA * threadIdx.y + threadIdx.x]; 
    Nds[threadIdx.y][threadIdx.x] = Nd[n + wA * threadIdx.y + threadIdx.x]; 
    __syncthreads(); 

    // keep track of the running sum  
    for (int k = 0; k < blockD; k++) 
     Pvalue += Mds[threadIdx.y][k] * Nds[k][threadIdx.x]; 
    __syncthreads(); 
    } 

    // write back to the global memory 
    int p = hB * blockD * blockIdx.y + blockD * blockIdx.x; 
    Pd[p + hB * threadIdx.y + threadIdx.x] = Pvalue; 
    __syncthreads(); 

    MaxFunction(Pd, max); 

} 

Главный код:

#include<stdio.h> 
#include "cuda.h" 
#include<stdlib.h> 

#define blockD 32 


const int wA = 128; 
const int hA = 1024; 

const int wB = 128; 
const int hB = wA; 

main(void){ 

    void MatrixMultiplication(float *, float *, float *, float *); 

    int size_A = wA * hA * sizeof(float); 
    int size_B = wB * hB * sizeof(float); 
    int size_C = wB * hA * sizeof(float); 
    int size_max = 2 * wB * sizeof(float); 
    float *M, *N, *P, *C; 


    // allocate memory on the CPU 
    M = (float*)malloc(size_A); 
    N = (float*)malloc(size_B); 
    P = (float*)malloc(size_max); 
    C = (float*)malloc(size_C); 

    // initialize the matrices 
    for (int y=0; y < hA; y++) { 
     for (int x=0; x < wA; x++){ 
      M[y*wA + x] = x; 
     } 
    } 

    for (int y=0; y<hB; y++) { 
     for (int x=0; x<wB; x++){ 
      N[y*wB + x] = x; 
     } 
    } 

    MatrixMultiplication(M, N, P, C); 

    //Write 
    FILE *f1; 
    int i, j; 
    f1 = fopen("max_val.txt","w"); 
    for(i=0; i < (wB * 2); i+=2){ 
    fprintf(f1,"%d\t%d\n",int(P[i]),int(P[i+1])); 
    } 
    fclose(f1); 

    f1 = fopen("Prod_mat.txt","w"); 
    for(i=0; i < 2; i++){ 
    for(j=0; j < wB; j++){ 
     fprintf(f1,"%d\t",int(C[i*wB + j])); 
    } 
    fprintf(f1,"\n"); 
    } 
    fclose(f1); 

    free(M); 
    free(N); 
    free(P); 
      free(C); 

    cudaDeviceReset(); 
    return 0; 
} 


void MatrixMultiplication(float *M, float *N, float *P, float *C) { 

    int size_A = wA * hA * sizeof(float); 
    int size_B = wB * hB * sizeof(float); 
    int size_C = wB * hA * sizeof(float); 
    int size_max = 2 * wB * sizeof(float); 
    float *Md, *Nd, *Pd, *max; 

    // allocate memory on the GPU 
    cudaMalloc((void**)&Md, size_A); 
    cudaMalloc((void**)&Nd, size_B); 
    cudaMalloc((void**)&Pd, size_C); 
    cudaMalloc((void**)&max, size_max); 

    // transfer M and N to device memory 
    cudaMemcpy(Md, M, size_A, cudaMemcpyHostToDevice); 
    cudaMemcpy(Nd, N, size_B, cudaMemcpyHostToDevice); 

    // kernel invocation code 
    dim3 dimBlock(blockD, blockD); 
    dim3 dimGrid(wA/blockD, hB/blockD); 

    //Execute Kernel 
    MatrixMulKernel<<<dimGrid, dimBlock>>>(Md, Nd, Pd, max); 

    // transfer P from device  
    cudaMemcpy(P, max, size_max, cudaMemcpyDeviceToHost); 
    cudaMemcpy(C, Pd, size_C, cudaMemcpyDeviceToHost); 

    cudaFree(Md); 
    cudaFree(Nd); 
    cudaFree(Pd); 
    cudaFree(max); 
} 

Результат умножения матриц отлично (проверено с помощью Matlab), но я не могу получить значения макс и их соответствующий индекс. Я был бы признателен, если кто-нибудь может любезно указать на то, что я делаю неправильно. Максимальная переменная имеет только мусор, когда я запускаю вышеуказанный код.

Answer 1

По-видимому, вы пытаетесь найти максимальное значение в каждом столбце, а также смещение к этому значению.

Но все ваши темы в y стучат в одном месте для максимального значения (max[x*2 + 0]). Это не рекомендуется, так как нет возможности разобраться в состоянии гонки. Вы должны использовать атомные операции или другие методы (например, сокращение), чтобы обрабатывать несколько потоков, обновляя одно максимальное значение таким образом.

Поскольку вам необходимо обновить два значения атомарно (максимальное значение и его местоположение), не просто заменить обычный доступ с помощью standard atomic function. Однако, поскольку вы имеете дело с двумя 32-битными смежными величинами, вы можете быть заинтересованы в моем ответе here.

Кстати, я думаю, что нативная матрица Matlab, умноженная на gpuArray, должна быть быстрее, чем любой код умножения матрицы, который вы пишете. Но для этого потребуется Parallel Compute Toolbox.