我看过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.xblockIdx.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.xblockIdx.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.xblockIdx.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.xblockIdx.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退散，风轻云淡”成就~
祝您好运！