关于使用流传输

我看过cuda 编程指南,它说:要使用多个流来传输数据,那么必须是page-locked memory上面的数据,但是又说这个memory的容量比较小,尽量少使用。到底这个容量是多大呢?我最近在用cuda,需要传输20MB的数据,数据只分配到global memory,但是数据传输到global memory就占了很大时间,很不划算啊。。。求指点。

LZ您好:

1:page-locked memory是host端的页锁定内存,和您cudaMalloc的device端的内存是两个概念。

2:按照现在机器的配置,您分配20MB的page-locked memory就和玩一样,您可以大胆地多分配一些无妨。

3:如果您的程序数据传输时间远多于计算时间,那么很可能是不适合在GPU上跑的,您需要先评估一下,再做决定。

祝您好运~

谢谢ICE…我试试

我还想问个有关constant memory的问题。是不是这样定义:
constant float A;
global void function()
{
}
int main
{

float H_a[10];
int s=10;
for(int i=0;i<s;i++)H_a[i]=0.0;
cudamalloc((void
*)&A,sizeof(float)*s);
cudaMemcpyTosymbal(A,h_A,sizeof(float)*s,cudaMemcpyHostToDevice);
//是使用cudaMemcpyTosymbal 去赋值给constant memory上面的变量的吧?
}

LZ您好:

您的代码虽然可能不报错,但是这样用是有问题的,您无法得到正确的结果。

您可以将“constant float *A;”改为“constant float A[10];”,并去除cudaMalloc,如此即可。

您原来的用法中,指针A在host端只能用ToSymbal系列函数调用,而不能用普通的需要一个存储在本地的指针作为参数的函数使用。

因为此时指针A是定义在device端的全局变量,而host端使用了特殊处理,保留了一个用于标记这个device端全局变量的同名变量,这个标记变量并不是实际的那个变量,也只能用在ToSymbal系列函数中。

ToSymbal系列函数会使用一套特殊处理的方法找到device端的对应变量。常规的函数只能直接使用host端的变量,因而是直接使用了该标记变量,这样并不能正确运作。

大致这样,供您参考。

祝您编码顺利~

以及,按照我5#给出的修改,您的A[10]是全部都标记为__constant__,他们全部会被constant cache缓冲。

而您4#的写法,只有指针A标记为__constant__,并被缓冲于constant cache中,cudaMalloc申请的空间则是常规的global memory(虽然4#写法实际无法正确申请)。
如果您要保持这种只把指针放在constant cache中的做法,那么可以如下改写:
constant float* A;
float *AA;
cudaMalloc(&AA, 10 * sizeof(float)):
cudaMemcpy(AA, H_A, 10 * sizeof(float), cudaMemcpyHostToDevice);
cudaMemcpyToSymbol(A, AA, sizeof(float *), cudaMemcpyHostToDevice);

请注意,此时只有指针A在constant memory中可以享受constant cache,10个元素依然是在global memory中的,并不享受constant cache的福利。

请您根据实际意图做出选择。

祝您好运~

constant_ float A[10];
global void function()
{
}
int main
{

float H_a[10];
int s=10;
for(int i=0;i<s;i++)H_a=0.0;
cudaMemcpyTosymbal(A,h_A,sizeof(float)*s,cudaMemcpyHostToDevice);
//是这样吗?
}

LZ您好,是这样的,这即是5#的建议,请您尝试即可。

祝您好运~

include “cuda_runtime.h”
include “device_launch_parameters.h”
include <stdio.h>
include <time.h>
include <stdlib.h>
define blocksize 16
constant float d_A[32];
constant float d_B[32];
global void addKernel(float *c, const float *a, const float b,const int num)
{
int i = threadIdx.x+blockDim.x
blockIdx.x;
if(i<num)
c[i] = a[i] + b[i];
}

int main()
{
float *A,*B,*C;

float *d_C;
int ss=32;

A=(float*)malloc(sizeof(float)ss);
B=(float
)malloc(sizeof(float)ss);
C=(float
)malloc(sizeof(float)*ss);
for(int i=0;i<ss;i++)
{
A[i]=1.0;
B[i]=2.0;
}

cudaMalloc((void**)&d_C,sizeof(float)*ss);

cudaMemcpyToSymbol(d_A,A,sizeof(float)*ss,cudaMemcpyHostToDevice);
cudaMemcpyToSymbol(d_B,B,sizeof(float)*ss,cudaMemcpyHostToDevice);
dim3 blocknum=(blocksize);
dim3 gridnum=((ss-1)/blocksize+1);
clock_t start=clock();
addKernel<<<gridnum,blocknum>>>(d_C,d_A,d_B,ss);

cudaDeviceSynchronize();
cudaMemcpy(C,d_C,sizeof(float)*ss,cudaMemcpyDeviceToHost);

clock_t end=clock();
for (int i=0;i<ss;i++)
{
	printf("%f\n",C[i]);
} 
printf("%f\n");
printf("time is %f\n",(double)(end-start)/CLOCKS_PER_SEC);

cudaFree(d_C);
free(A);
free(B);
free(C);
return 0;

}
这是我测试的代码,运行没异常。。但是结果都是0.000,不知道为什么?帮我看看吧,谢谢 了·(^o^)/~

LZ您好,看了您的代码

1:您没有理解__constant__(以及类似的__device__)定义变量的实际含义。因而混用了__constant__定义的数组和host端定义指针并指向device端内容的两种用法。

在5#已经指出,constant 定义的A,在host端直接访问,并非是原本的A,而是host端的一个标记,只能用于ToSymbol系列函数,而您却将这样的标记作为kernel函数的参数,这样必然是不行的。

2:如果您打算保留__constant__ float d_A[32];constant float d_B[32];这样的定义,那么,请将您的kernel函数修改为:

global void addKernel(float c, const int num)
{
int i = threadIdx.x+blockDim.x
blockIdx.x;
if(i<num)
c [ i ]= d_A[ i ] + d_B[ i ];
}

并将host端调用kernel代码改为:
addKernel<<<gridnum,blocknum>>>(d_C,ss);

即可。

3:如果您打算使用指针传递参数的形式,那么一般可以

float* d_A,d_B;
cudaMalloc(&d_A,…);
cudaMalloc(&d_B,…);

cudaMemcpy(d_A,…);
cudaMemcpy(d_B,…);

addKernel<<<…>>>(d_C,d_A,d_B,ss);

以及您的kernel可以不改,仍为:
global void addKernel(float *c, const float *a, const float b,const int num)
{
int i = threadIdx.x+blockDim.x
blockIdx.x;
if(i<num)
c[ i ] = a[ i ] + b[ i ];
}

这样写是可行的。

但是请不要把两种写法混合起来。

4:以及,您会发现3:中的写法没有使用constant cache,是的,因为您这样使用__constant__,会导致一个warp内部各线程每个线程访问的都是不同的位置,而constant cache需要在warp内各线程访问一致或基本一致的时候才能取得较好的效果。您的做法会导致严重的性能问题。

5:以及,如果您坚持使用9#中的d_A,d_B的定义,以及坚持使用9#的kernel写法,那么请使用cudaGetSymbolAddress()这个函数重新取得d_A和d_B的实际device端地址,并赋值给本地的float*d_A_h,d_B_h,然后使用d_A_h和d_B_h作为kernel 的参数即可。

6:以及,如果您希望使用2:中的device端全局变量的写法,并且此时放弃使用__constant__,避免不合适地使用constant cache以降低性能,那么您可以简单将2:中定义d_A和d_B时的__constant__ 改为__device__,即可。

大致这样,祝您编码顺利~

谢谢你。。不知道为何,我计算机跑出来的结果还是 0.00, 都按照你说的去写了。。无奈了···哎· ·

您确定么?发代码吧。别空说折腾ICE, 最后你还写的不是ICE说的。

LZ您好,您的程序改的如何了?本帖涉及内容都是常规内容的,一般不会出什么幺蛾子的。

祝您好运~

我使用的是GF405M,比较老的显卡,1.2X

include “cuda_runtime.h”
include “device_launch_parameters.h”
include <stdio.h>
include <time.h>
include <stdlib.h>
define blocksize 16
constant float d_A[32];
constant float d_B[32];
global void addKernel(float c,const int num)
{
int tx=blockDim.x
blockIdx.x+threadIdx.x;
if(tx<num)
c[tx] =d_A[tx]+d_B[tx];
}

int main()
{
float *A,*B,*C;
float d_A,d_B;
float d_C;
const int ss=32;
A=(float
)malloc(sizeof(float)ss);
B=(float
)malloc(sizeof(float)ss);
C=(float
)malloc(sizeof(float)ss);
for(int i=0;i< ss;i++)
{
A[i]=1.0;
B[i]=2.0;
}
cudaDeviceProp deviceProp;
cudaMalloc((void
)&d_C,sizeof(float)
ss);

cudaMemcpyToSymbol(d_A,A,sizeof(float)* ss,cudaMemcpyHostToDevice);
cudaMemcpyToSymbol(d_B,B,sizeof(float)* ss,cudaMemcpyHostToDevice);
cudaThreadSynchronize();

dim3 blocknum=(blocksize);
dim3 gridnum=(( ss-1)/blocksize+1);
clock_t start=clock();

addKernel<<<gridnum,blocknum>>>(d_C,ss);

cudaMemcpy(C,d_C,sizeof(float)*ss,cudaMemcpyDeviceToHost);

clock_t end=clock();
for (int i=0;i< ss;i++)
{
	printf("%f\n",C[i]);
} 
printf("%f\n");
printf("time is %f\n",(double)(end-start)/CLOCKS_PER_SEC);

cudaFree(d_C);
free(A);
free(B);
free(C);
return 0;

}

LZ您好:

请把您main()函数中的“float *d_A,*d_B;”这一行拿掉。

以及,您的计算能力1.2的卡并无支持问题。

祝您好运~

我把这行注释掉了,还是不对的。你能运行结果是对的?假如可以,那我想是我的编译器问题了吧 cuda5.0+win7+vs2008.

LZ您好:

之前建议您拿掉main()函数里面的float* d_A,*d_B是因为,他们会掩盖您定义的全局数组d_A[32]和d_B[32]。

此外,刚才发现,您的cudaMemcpyToSymbol()的参数使用不正确,您写“cudaMemcpyHostToDevice”这个参数实际上是在该函数的“offset”参数的位置,而并不在“cudaMemcpyKind”这个参数的位置。
所以,请您参照手册的说明,修正这一点。

因为我不常用ToSymbol类的函数,这个问题之前未能及时发现,深表歉意。

祝您调试顺利~

嗯嗯,你说对了。程序对了·谢谢你O(∩_∩)O

善哉,恭喜LZ解开“BUG退散,风轻云淡”成就~

祝您好运!