Slow local memory, feigned constant memory. coalesced? global?

[tonhead]

Dear Cygnus X1,

Thank you for your helpful comments.

I will follow your path.

============================

All of your code examples are completly serial. All threads perform exactly the same thing, which is kinda pointless. Are you aware of that?

1. They are serial, they are doing almost the same work because it is Monte-Carlo algorithm implementation. But the code I posted here is not the whole program, there are some other calculations going on there.

============================

Secondly:

__constant__ double* constant_a;

This means that a pointer to double resides in constant memory, but the actual data is not! As a result your line

volatile_a[i][j] = constant_a[i * 16 + j];

will be compiled to something like this:

offset1:=i * 16 + j;

constant_a_base:=constant_memory[constant_a];

var:=global_memory[constant_a_base+offset1];

local_memory[volatile_a+offset1]:=var;

If you want a whole array to be in constant memory, write this instead:

__constant__ constant_a;

2. That makes sense! Same as const in C++. So in order to place the data in constant I do

__constant__ double constant_a[16][16];

============================

Thirdly:

error_here = cudaMemcpy(constant_a, host_constant_a, size_constant_a, cudaMemcpyHostToDevice);

constant_a (as it is declared so far) has a value which resides on GPU, thus cannot be accessed directly on the host side. I am surprised the compiler did not complain about this thing. Unless you use cudaemu, you should read some garbage instead. So you move host_constant_a data to some random position in GPU memory.

Similarly &constant_a should raise a compile error. You use cudaGetSymbolAddress (or something like that) to get address of device (or constant?) variables. Alternatively you may want to use cudaCopyToSymbol (or something like that) function to store data into constant memory.

3. This was taken from the NVIDIA CUDA Programming Guide, where they do

float* d_A; 

	cudaMalloc((void**)&d_A, size); 

	cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);

and it works fine for me. Sure, in the case if we have data in the constant memory, we want to get an address of it in the first place. Thank you for mentioning that.

Constant variables are like global variables, they are allocated upon declaration so no need to call for cudaMalloc on them.

It’s the first time I hear so.

============================

Fourthly:

double volatile_a[16][16];

Declares an array for each thread. Since it is big, it cannot fit into 16 (or 32) of 32-bit registers, it is stored in local memory. Physically local memory is a global memory. As a result all you do in that loop is to copy data from one place in global memory to another place in global memory!

Sure it is in local :) I think it is no doubt in the local memory, also because indexing is not static, loops are not unrolled.

Yes, all I do is copying data from constant memory space to global local memory, and then I do some operations on that local data (solve the system, to be exact, function is given in my second post here).