【问题标题】:transpose different matrices in parallel并行转置不同的矩阵
【发布时间】:2019-09-23 21:21:45
【问题描述】:

我有 3 个不同大小的矩阵,并希望将它们并行转置。

首先我使用 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;
}

【问题讨论】:

  • NM 在编译时总是已知的吗?您可以将内核设为模板并使用固定大小的数组。
  • 您将A 更改为指向其他位置,然后调用delete。即使在主机代码中也是一个错误。
  • 我不熟悉 cublas,但是这些函数真的是要从内核调用的吗?否则你可能会查看cublasSetStream
  • 可以从设备调用cublas函数。我认为重点是如何从数组中获取矩阵的相应数据。并且设备中的 malloc A 和 C 可能是错误的?
  • 看来你是对的,有可能:docs.nvidia.com/cuda/cuda-c-programming-guide/…

标签: c++ cuda


【解决方案1】:

您的代码中有许多错误。我的描述中可能会遗漏一些内容。

  1. 请注意,此 cublas-in-device-code 功能在较新的 CUDA 版本中不再可用。
  2. 传递给设备代码的每个指针都需要使用cudaMalloc 进行分配。您已经完成了cudaMalloc 的一些指示,但不是全部。
  3. 您对指针和指针数组感到困惑。我将无法为您解决所有这些问题。您的内核设计确实不需要使用指针数组的复杂性。所以我已经删除了所有这些。
  4. 在 CUDA 动态并行 (CDP) 中,无法将指向本地地址空间的指针传递给子内核。您不能在本地地址空间中使用 alpha 和 beta,并将指向它们的指针传递给 CDP 中的 CUBLAS。
  5. 要进行纯转置,请研究 the CUBLAS Sgeam documentation 以了解推荐使用的参数。

我相信我还修复了其他问题。请研究这个例子:

$ 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
$

【讨论】:

  • 非常感谢!我成功运行它。我仍然困惑只有常量值不需要传递给设备?
  • 您需要对 C 编程语言中按值传递的含义有一个非常清晰的理解。所有内核参数都是按值传递的。取消引用任何指针都必须产生正确的设备地址。
  • 对不起,错误的表达方式。我的意思是只有常量值可以直接传递给设备,而其他人需要使用 cudaMalloc 分配?
  • 是的,我也是这个意思。所有内核参数都是按值传递的。对于标量,按值传递就可以了。对于指针数量,指针value(即地址)将按值传递。当您取消引用该指针(即使用地址检索值)时,address 最好指向设备内存。因此任何时候你传递一个指向设备的指针,该指针最好指向设备内存。这是一个基本的 CUDA 概念。您不能取消引用主机代码中的设备指针或设备代码中的主机指针。
  • 但是您没有取消引用标量。所以按值传递在那里工作得很好,你可以直接在设备代码中使用数量。您确实需要更好地理解指针和按值传递,这些都是 C 编程概念,不是 CUDA 独有的或特定于 CUDA。
猜你喜欢
  • 2013-05-27
  • 1970-01-01
  • 2017-11-23
  • 1970-01-01
  • 1970-01-01
  • 1970-01-01
  • 2021-07-05
  • 1970-01-01
相关资源
最近更新 更多