sorry, one more question, does anyone has an example to pass kernel arguments as constant memory? thanks
-
copy data from host(or device) memory to const memory (using cudaMemcpyToSymbol()).
-
you not need to pass these arguments to kernel
-
using arguments as public variable.
for example,
[b]device constant float constArray[n];
//copy from host to const memory
cudaMemcpyToSymbol(constArray, hostArray, n * sizeof(float), cudaMemcpyHostToDevice);
//using const memory inside kernel
global void kernel(/some arguments/)
{
int tid = threadIdx.x;
//access data allocated in const memory
float temp = constArray[tid];
//do some thing more
}[/b]
thank you for the example. I guess these can only be defined as the global variables.
ask a follow up question, in the “Best Practics Guide”, it says “The constant memory space is cached… reading from the constant cache is as fast as reading from a register …”. Where are these constant memory variables cached? in the shared memory? register? or a separate high-speed memory as fast as registers?
If they are cached in the shared memory, then moving kernel argument from shared mem to constant mem will not save anything, so, it must not, right?