Uncoalesced on matrix by vector multiplication

[danielfleischman]

Hi,

I am making an application which need to multiply a (row) vector by a matrix (y = x*A). So I’ve done this code:

__global__ void calc(float* mat, float* in, float* out, int tam) {

  int ix = blockIdx.x*blockDim.x + threadIdx.x;

  float ans = 0.;

  int j = 0;

  for (int i = ix; i < tam*tam; i+=tam) {

	ans += mat[i]*in[j];

	j++;

  }

  out[ix] = ans;

}

It is correct, I’ve tested it against several matrices/vectors and it is all fine.

But the problem is that when I run the cudaprof with a small example (64x64) and I get the following result: “gld coalesced = 128” and “gld uncoalesced = 2048”. I don’t know where does these uncoalesced access come from.

Can someone help me?

Thank you,

Daniel.

PS.: The “in” is declared as constant

[avidday]

You will get uncoalesced access on 1.1 hardware any time threads in a half-warp don’t sequentially access contiguous 32/64/128 byte blocks of global memory which are aligned on 32/64/128 byte boundaries.

So (by my rather inexpert eye) your loads of in will never coalesce (every thread in the half warp reads the same value at each loop iteration), and your loads of mat would only coalesce when the stride is aligned to 64 byte boundaries.

[danielfleischman]

You will get uncoalesced access on 1.1 hardware any time threads in a half-warp don’t sequentially access contiguous 32/64/128 byte blocks of global memory which are aligned on 32/64/128 byte boundaries.

So (by my rather inexpert eye) your loads of in will never coalesce (every thread in the half warp reads the same value at each loop iteration), and your loads of mat would only coalesce when the stride is aligned to 64 byte boundaries.

Thank you for your answer, but from the NVIDIA CUDA Programming Guide:

For all threads of a half-warp, reading from the constant cache is as fast as reading

from a register as long as all threads read the same address. The cost scales linearly

with the number of different addresses read by all threads. We recommend having

all threads of the entire warp read the same address as opposed to all threads within

each of its halves only, as future devices will require it for full speed read.

That’s what I tried to do. It seems like I need to better manage my cache… any suggestion?

Tjank you,

Daniel.

[avidday]

TThat’s what I tried to do. It seems like I need to better manage my cache… any suggestion?

Re-read the complete chapter of the user guide you quoted from, because you seem to have misunderstood it. Your arrays are loading and storing to global memory, which is uncached. Constant memory and constant memory caching have no bearing on your problem.

You can probably coalesce the loads by having each thread perform staged loads to shared memory in front of a synchronization barrier first, then read from shared memory inside the loop. There is only 16k of shared memory per block, so you will have to think carefully about block sizes and, (perhaps) decomposing the operation into a smaller set of sub calculations which can work inside the shared memory limit . There is a very good set of slides written by Mark Harris from NVIDIA from SC’08 which discuss memory coalescing strategies in detail. You might find useful to study.