最新帖子

页: [1] 2 3 ... 10
1
我这结果就不正确,设置分块矩阵那里蒙蔽了
2
#include <cuda_runtime.h>
#include <cublas_v2.h>
#include <iostream>
#include <stdio.h>

using namespace std;
#define IDX2F(i,j,ld) ((((j)-1)*(ld))+((i)-1))

#define IDX2C(i,j,ld) (((j)*(ld))+(i))
//ld是维度,i是行,j是列,cuBLAS使用的是列存储
//这个宏表示, 矩阵第i行第j列的元素位置在C语言中 数组存储的索引
//ld表示 矩阵的第一个维的元素个数,就是 矩阵的行数。
__global__
void show(float* ptr, int size)
{
    for (int i = 0; i < size; i++)
        printf("%f\n", ptr);
}


void print_matrix(int R, int C, float* A, const char* name)
{
    printf("%s = \n", name);
    for (int r = 0; r < R; ++r)
    {
        printf("[");
        for (int c = 0; c < C; ++c)
        {
            printf("%10.6f", A[c * R + r]);
        }
        printf("] \n");
    }
}


void print_matrix_(int R, int C, float* A, const char* name)
{
    printf("%s = \n", name);
    for (int r = 0; r < R; ++r)
    {
        printf("[");
        for (int c = 0; c < C; ++c)
        {
            printf("%10.6f", A[r * C + c]);
        }
        printf("]\n");
    }
}

int main()
{
    int M = 3; //行数 矩阵A的行,结果矩阵C的行数.A=[3,9]
    int N = 9; //列数 矩阵A的列,矩阵B的列   B=[3,9]
    int B = 3; //行数 矩阵B的行
    int K = 3; //列数,结果矩阵C的列数  C=[3,3]

    //分配主机矩阵并初始化
    float* a, * b, * c;
    cudaHostAlloc((void**)&a, sizeof(float) * M * N, cudaHostAllocDefault);
    cudaHostAlloc((void**)&b, sizeof(float) * B * N, cudaHostAllocDefault);
    cudaHostAlloc((void**)&c, sizeof(float) * M * K, cudaHostAllocDefault);

    for (int j = 0; j < N; j++) {
        for (int i = 0; i < M; i++) {
            a[IDX2C(i, j, M)] =  1.0;
        }
    }

    /*
    for (int i = 0; i < M; ++i) {
        for (int j = 0; j < N; ++j) {
            a[i * N + j] = 1;
        }
    }

   
    print_matrix_(M, N, a, "A");

   
    */

    //可视化矩阵
    print_matrix(M, N, a, "A");

    for (int j = 0; j < N; j++) {
        for (int i = 0; i < B; i++) {
            b[IDX2C(i, j, B)] = (float)(i * N + j + 1);
        }
    }

    /*
    for (int i = 0; i < B; ++i) {
        for (int j = 0; j < N; ++j) {
            b[i * N + j] = i * N + j + 1;
        }
    }
    */
    //print_matrix(B, N, b, "B");

    print_matrix(B, N, b, "B");
    //分配设备的数据
    float* d_a, * d_b, * d_c;
    cudaMalloc(&d_a, sizeof(float) * M * N);
    cudaMalloc(&d_b, sizeof(float) * B * N);
    cudaMalloc(&d_c, sizeof(float) * M * K);

    //Host->device
    cudaMemcpy(d_a, a, sizeof(float) * M * N, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, b, sizeof(float) * B * N, cudaMemcpyHostToDevice);
    cudaMemcpy(d_c, c, sizeof(float) * M * K, cudaMemcpyHostToDevice);
    //

    cublasHandle_t handle;
    cublasStatus_t ret;
    ret = cublasCreate(&handle);

    //矩阵分块计算
    float* a_array[9], * b_array[9];
    float* c_array[9];

    int r = 3;
    int l = 3;

    for (int i = 0; i < r; ++i) {
        for (int j = 0; j < l; ++j) {
            a_array[i * l + j] = d_a + i * 9 + j * r;
            b_array[i * l + j] = d_b + i * 9 + j * r;
            c_array[i * l + j] = d_c + i * 9 + j * r;
            //printf("%d\n",*a_array[i * l + j]);
        }

    }
    //print_matrix_(r, l, *(a_array), "a_array");
    const float** d_Marray, ** d_Narray;
    float** d_Parray;
    cudaMalloc((void**)&d_Marray, N * sizeof(float*));
    cudaMalloc((void**)&d_Narray, N * sizeof(float*));
    cudaMalloc((void**)&d_Parray, N * sizeof(float*));
    cudaMemcpy(d_Marray, a_array, N * sizeof(float*), cudaMemcpyHostToDevice);
    cudaMemcpy(d_Narray, b_array, N * sizeof(float*), cudaMemcpyHostToDevice);
    cudaMemcpy(d_Parray, c_array, N * sizeof(float*), cudaMemcpyHostToDevice);


    const float alpha = 1.0f;
    const float beta = 0.0f;

    //需要的是A矩阵的一个1行3列的矩阵乘以矩阵b的三行一列,
    int m = 1; //按列 m = 1
    int n = 3; //按列 n = 3
    int k = 1; //按列 k = 1

    int lda = 9;
    int ldb = 9;
    int ldc = 9;
    int batch = 9;
    //    矩阵OP(A)的维度是m×k
    //    矩阵OP(B)的维度是k×n
    //    矩阵C的维度是m×n
    //    运算为C = alpha * A * B + beta * C
    ret = cublasSgemmBatched(handle,
        CUBLAS_OP_N,
        CUBLAS_OP_N,
        k, m, n,
        &alpha,
        d_Narray, ldb,
        d_Marray, lda,
        &beta,
        d_Parray, ldc,
        batch);

    cublasDestroy(handle);
   
    if (ret == CUBLAS_STATUS_SUCCESS)
    {
        printf("sgemm success  %d, line(%d)\n", ret, __LINE__);
    }
   
    //show << <1, 1 >> > (c_array[0], 16);
    cudaMemcpy(c, d_c, sizeof(float) * M * K, cudaMemcpyDeviceToHost);

    print_matrix(M, K, c, "C = A x B");
    return 0;

}

3
对不起我表述的不是很清楚,就是这个意思QAQ
4
前辈您好!我的第一个问题是想问我的event的使用是否会和数据传输起到冲突,因为cudamemcpy也存在着隐式同步,我担心它与cudaEventSynchronize()起了冲突。我对event的使用:创建event,指定cudaEventBlockingSync,然后record,最后cudaEventSynchronize(),这个流程应该没有问题吧QAQ,感谢您的回复!

我看不懂你的前半部分。。。你可以在一个流里下发cudaMemcpyAsync(),然后后面立刻随着一个event的record(你的后半段),然后等待event的record()在此流中完成。---这是你的意思?
5
前辈您好!我的第一个问题是想问我的event的使用是否会和数据传输起到冲突,因为cudamemcpy也存在着隐式同步,我担心它与cudaEventSynchronize()起了冲突。我对event的使用:创建event,指定cudaEventBlockingSync,然后record,最后cudaEventSynchronize(),这个流程应该没有问题吧QAQ,感谢您的回复!
6
CUDA / Re: cuda 矩阵计算思路求教!!!
« 最后发表 作者 塬上的卡夫卡 六月 30, 2021, 05:46:02 pm »
(1)缓冲没有必要一次性全部缓冲完。你可以看一下手册自带的"矩阵乘法"的例子, 里面完全不是一次性缓冲完的(可以将你这个看成它的特例,想法是类似的)。
(2)如果你要缓冲M, 最坏64x64*16B(2个double), 这需要64KB, 一次性放入shared memory不可以,但是可以考虑分片,或者直接读取寄存器中, 这是可以足够放入的. 然后需要跨线程累加起来得到标量不过(直接用shuffle交换累加)。或者如果你适当的转置(和你的存储格式有关),再每次点乘,也可以不用考虑跨线程累加。

(3)从CPU端调用cublas, 和动态并行无关
(4)从设备端如果想调用,还有动态并行方面的顾虑,则现在不支持CUDA DP的显卡已经停止支持了(NV目前的Toolkit从3.5+开始支持,甚至下一个版本3.X将被彻底拿掉). 所以不用考虑这个原因。
对,听了你的建议,把cuda sample好好看了下,终于打开了思路,问题解决了,效率一下提高了不少,感谢!
7
所谓的线程busy polling是线程一旦调用设备(GPU)就会busy polling吗,亦或者是GPU触发同步语句包括不限于cudamecpy,才会busy polling?QAQ

你的问题的第一部分,python中的CPU(), 这个我不懂。在最初的时候说过了,欢迎其他使用python/pytorch的用户回答。

关于你的问题的第二部分,何时会spin,根据NV的官方说法:
“当前系统上的CPU核心数量,超过GPU卡的个数的时候,则在CPU<->GPU的时候,会进行忙等“。而没有提到cudaMemcpy()的时候的具体策略。

我建议总是启用使用同步对象的等待(立刻释放时间片)。你可以总是在设备级别或者Event级别启用它。
8
所谓的线程busy polling是线程一旦调用设备(GPU)就会busy polling吗,亦或者是GPU触发同步语句包括不限于cudamecpy,才会busy polling?QAQ
9
您好!设定cudaEventBlockingSync标志,这个阻塞host方法粒度更小,应该是更好的选择。但是我尝试过:pytorch封装了Event和Stream对象,但是我尝试了调用,并不能解决,我已在pytorch官网提了issue(详情见:https://github.com/pytorch/pytorch/issues/60541),并没有得到回复。关于ctypes调用cudart64.dll里面的函数,我认为他应该是调用失败了,cudaSetDeviceFlags之后调用cudaGetDeviceFlags返回的结果并不正确。我现在尝试先用C++调用cudaSetDeviceFlags然后重新封装为dll,再到python中调用。然后我还有一个疑问万望解答,pytorch中CPU()方法用于将GPU数据复制回CPU,我不确定他是异步还是同步。如果使用cudaEventSynchronize()应该在CPU()之前还是之后呢。代码可以在issue中看到。感谢您的回复!
10
https://bbs.gpuworld.cn/index.php?topic=58836.0,如这篇文章所说,我想设定BlockingSync的设备标志,来阻塞host线程以获取更高的性能,否则CPU总是处于100%的自旋轮询。但是Pytorch没有这个选项,所以我考虑通过用ctypes调用cudaSetDeviceFlags,但是应该是失败,总之没有效果。pyqt的线程依然是busy polling的状态。这个问题困扰了我很久,有什么解决办法吗?不胜感激!盼回复。

我不懂python/pytorch, 但是听你的描述像是在不正确的时机设定了标志(cudaDeviceScheduleBlockingSync需要在最开头就设定好),或者你的时机正确,但是被pytorch内部给重新设定了,这个时候可能并没有什么好的办法。

一个可能的解决方案是,你额外使用ctypes对event对象设定cudaEventBlockingSync, 这个是独立的,而且可以在每次同步都选择不同的方式(例如对预计的短kernel选择spin, 或者普通kernel选择阻塞同步),而且这个可以在任意中途创建event并使用。

也欢迎你给出pytorch的原生解决方案(假设你修改了它的源代码),并发到论坛,方便其他的论坛的兄弟姐妹们。
页: [1] 2 3 ... 10