транспонировать разные матрицы параллельно

Question

Я имею как 3 разных размера матриц и хочу транспонировать их параллельно.

Сначала я помещаю их в 2D-массив с помощью malloc, а затем использую cudaMalloc для передачи массива с хоста (h_B) на устройство (d_B).

Использование threadIdx для поиска каждого адреса матрицы в массиве. Функция cublas используется. Вот мой код.

Код может быть скомпилирован, но я не могу получить результат. Кажется, что в глобальной функции float *A = new float[m*n] не очень хороший способ.

У кого-нибудь есть идеи на этот счет? Большое спасибо!

/* Includes, system */
#include <stdio.h>
#include <stdlib.h>
#include<iostream>

/* Includes, cuda */
#include <cuda_runtime.h>
#include <cublas_v2.h>

/* Includes, cuda helper functions */
#include <helper_cuda.h>

__global__ void transposeCublasSgeam(int *M_A, int *N_A, float *ptrA, float *ptrC, const int N, int *address)
{
    cublasHandle_t cnpHandle;
    cublasStatus_t status = cublasCreate(&cnpHandle);

    if (status != CUBLAS_STATUS_SUCCESS)
    {
        return;
    }

    const float d_alpha = 1.0f;
    const float d_beta = 0.0f;
    int idx = threadIdx.x;
    if(idx<N){
      int m = M_A[idx]; //A_row
      int n = N_A[idx]; //A_col
      float *A = new float[m*n]; 
      float *C = new float[m*n];
      A = ptrA+address[idx];
      C = ptrC+address[idx];
      cublasSgeam(cnpHandle, CUBLAS_OP_T, CUBLAS_OP_T, m, n, &d_alpha, (const float*)A, n, &d_beta, (const float *)A, n, C, m);
      delete[] A;
      delete[] C;
    }
     cublasDestroy(cnpHandle);

}

int main()
{

    const int N = 3;
    int M_B[N] = { 2,3,2 }; //row number of matrices 
    int N_B[N] = { 3,2,4 }; //col number of matrices 

    float a[6] = { 1,2,3,
                  4,5,6 };
    float b[6] = { 1,2,
                 3,4,
                 5,6};
    float c[8] = { 1,2,3,1,
                  2,3,4,5 };

    float **h_B = (float**)malloc(N * sizeof(float*));
    float **h_BT = (float**)malloc(N * sizeof(float*));

    h_B[0] = a, h_BT[0] = a;
    h_B[1] = b, h_BT[1] = b;
    h_B[2] = c, h_BT[2] = c;

    int NUM_B = 20; // total number of elements
    int address[] = {0,6,12}; 

    float *d_B, *d_BT;
    checkCudaErrors(cudaMalloc((void **)&d_B, NUM_B * sizeof(float)));
    checkCudaErrors(cudaMalloc((void **)&d_BT, NUM_B * sizeof(float)));
    checkCudaErrors(cudaMemcpy(d_B, h_B, NUM_B * sizeof(float), cudaMemcpyHostToDevice));
    checkCudaErrors(cudaMemcpy(d_BT, h_BT, NUM_B * sizeof(float), cudaMemcpyHostToDevice));

    transposeCublasSgeam<<<1,N>>>(M_B, N_B, d_B,d_BT, N,address); 

    checkCudaErrors(cudaMemcpy(h_BT, d_BT, NUM_B * sizeof(float), cudaMemcpyDeviceToHost));

    cudaFree(d_B);
    cudaFree(d_BT);
    delete[] h_B;
    delete[] h_BT;

    return 0;
}

Answer 1

В вашем коде был ряд ошибок. Я, вероятно, пропущу некоторые из них в своем описании.

1. Обратите внимание, что эта функция кода в устройстве больше не доступна в более новых версиях CUDA.
2. Каждый указатель, который передается коду устройства, должен быть выделен с cudaMalloc. Вы сделали cudaMalloc за несколько указателей, но не все из них.
3. Вы запутались в указателях и массивах указателей. Я не смогу разобраться со всем этим для вас. Ваш дизайн ядра действительно не нуждается в сложности использования массивов указателей. Так что я все это удалил.
4. В динамическом параллелизме CUDA (CDP) указатели на локальное адресное пространство не могут передаваться дочерним ядрам. Вы не можете использовать альфа и бета в локальном адресном пространстве и передавать указатели на них в CUBLAS в CDP.
5. Чтобы выполнить чистую транспонирование, изучите документацию CUBLAS Sgeam для рекомендуемых параметров для использования.

Я считаю, что были другие вещи, которые я исправил. Пожалуйста, изучите этот пример:

$ cat t1433.cu
/* Includes, system */
#include <stdio.h>
#include <stdlib.h>
#include<iostream>

/* Includes, cuda */
#include <cuda_runtime.h>
#include <cublas_v2.h>

/* Includes, cuda helper functions */
#include <helper_cuda.h>

__global__ void transposeCublasSgeam(int *M_A, int *N_A, float *ptrA, float *ptrC, const int N, int *address)
{
    cublasHandle_t cnpHandle;
    cublasStatus_t status = cublasCreate(&cnpHandle);

    if (status != CUBLAS_STATUS_SUCCESS)
    {
        printf("thread: %d, error1: %d\n", threadIdx.x, (int)status);
        return;
    }

    float *d_alpha =  new float; // a pointer to device-heap, not local memory
    *d_alpha = 1.0f;
    float *d_beta  = new float;
    *d_beta = 0.0f;
    int idx = threadIdx.x;
    if(idx<N){
      int m = M_A[idx]; //A_row
      int n = N_A[idx]; //A_col
      status = cublasSgeam(cnpHandle, CUBLAS_OP_T, CUBLAS_OP_N, m, n, d_alpha, ptrA+address[idx], n, d_beta, ptrC+address[idx], m, ptrC+address[idx], m);
      if (status != CUBLAS_STATUS_SUCCESS)
      {
        printf("thread: %d, error2: %d\n", threadIdx.x, (int)status);
        return;
      }
    }
     cublasDestroy(cnpHandle);

}

int main()
{

    const int N = 3;
    int M_B[N] = { 2,3,2 }; //row number of matrices
    int N_B[N] = { 3,2,4 }; //col number of matrices

    float a[6] = { 1,2,3,
                  4,5,6 };
    float b[6] = { 1,2,
                 3,4,
                 5,6};
    float c[8] = { 1,2,3,1,
                  2,3,4,5 };
    float *h_Bdata  = (float *)malloc(sizeof(a)+sizeof(b)+sizeof(c));
    float *h_BTdata = (float *)malloc(sizeof(a)+sizeof(b)+sizeof(c));
    memcpy(h_Bdata, a, sizeof(a));
    memcpy(h_Bdata+(sizeof(a)/sizeof(a[0])), b, sizeof(b));
    memcpy(h_Bdata+(sizeof(a)/sizeof(a[0]))+(sizeof(b)/sizeof(b[0])), c, sizeof(c));

    int NUM_B = 20; // total number of elements
    int address[] = {0,6,12};
    int *d_address;
    cudaMalloc(&d_address, sizeof(address));
    cudaMemcpy(d_address, address, sizeof(address), cudaMemcpyHostToDevice);
    int *d_M_B, *d_N_B;
    cudaMalloc(&d_M_B, sizeof(M_B));
    cudaMalloc(&d_N_B, sizeof(N_B));
    cudaMemcpy(d_M_B, M_B, sizeof(M_B), cudaMemcpyHostToDevice);
    cudaMemcpy(d_N_B, N_B, sizeof(N_B), cudaMemcpyHostToDevice);
    float *d_B, *d_BT;
    checkCudaErrors(cudaMalloc((void **)&d_B, NUM_B * sizeof(float)));
    checkCudaErrors(cudaMalloc((void **)&d_BT, NUM_B * sizeof(float)));
    checkCudaErrors(cudaMemcpy(d_B, h_Bdata, NUM_B * sizeof(float), cudaMemcpyHostToDevice));

    transposeCublasSgeam<<<1,N>>>(d_M_B, d_N_B, d_B,d_BT, N,d_address);

    checkCudaErrors(cudaMemcpy(h_BTdata, d_BT, NUM_B * sizeof(float), cudaMemcpyDeviceToHost));
    std::cout << "B , BT" << std::endl;
    for (int i = 0; i < NUM_B; i++){
      std::cout << h_Bdata[i]  << " , " << h_BTdata[i] << std::endl;}
    cudaFree(d_B);
    cudaFree(d_BT);

    return 0;
}
$ /usr/local/cuda-8.0/bin/nvcc -I/usr/local/cuda-8.0/samples/common/inc  t1433.cu -rdc=true -lcublas_device -lcudadevrt -arch=sm_35 -o t1433
$ LD_LIBRARY_PATH=/usr/local/cuda-8.0/lib64 CUDA_VISIBLE_DEVICES="3" cuda-memcheck ./t1433
========= CUDA-MEMCHECK
B , BT
1 , 1
2 , 4
3 , 2
4 , 5
5 , 3
6 , 6
1 , 1
2 , 3
3 , 5
4 , 2
5 , 4
6 , 6
1 , 1
2 , 2
3 , 2
1 , 3
2 , 3
3 , 4
4 , 1
5 , 5
========= ERROR SUMMARY: 0 errors
$