Dnia 2010-12-13, pon o godzinie 19:13 -0800, reckoner pisze: > thank you. These comments are very helpful. > > I don't understand what you mean by the following: > > "You are also using full expression to calculate index: > const int i = blockDim.x*blockIdx.x + threadIdx.x; > while you are using only one block - this makes your code > harder to analyse." > > How else should I be calculating the index? >
I meant that since you are using only one block,
it is enough to write:
const int i = threadIdx.x;
instead of full expression.
Using
const int i = blockDim.x*blockIdx.x + threadIdx.x;
suggests that you are using many blocks.
This is not an error, but I think it it important to make
sure that code really says what do we mean.
For example look at code I have just sent, wrapping CURAND library:
__global__ void uniform_int(curandStateSobol32 *s, int *d, int n) {
const int tidx = threadIdx.x;
for (int idx = tidx; idx < n; idx += blockDim.x) {
d[idx] = curand(&s[tidx]);
}
}
self.uniform_int.prepare("PPi", block=(generator_count, 1, 1))
self.uniform_int.prepared_async_call((1, 1), stream, self.state, data,
input_size)
I am using only one block so I have tidx = threadIdx.x
At the same time I want my thread to take care of many values
(that would normally be served by separate blocks), so I have
put idx += blockDim.x into a loop.
Omission of blockDim.x*blockIdx.x from expression computing tidx
was done on purpose, to point to someone reading code that I am using
threads from one block.
Now I have started wondering whether not to rewrite loop to say:
for (int idx = threadIdx.x; idx < n; idx += blockDim.x) {
but I am not yet sure whether to do it or not.
Regards.
> Thanks again.
>
> On 12/13/2010 3:49 PM, Tomasz Rybak wrote:
> > You are also using full expression to calculate index:
> > const int i = blockDim.x*blockIdx.x + threadIdx.x;
> > while you are using only one block - this makes your code
> > harder to analyse.
--
Tomasz Rybak <bogom...@post.pl> GPG/PGP key ID: 2AD5 9860
Fingerprint A481 824E 7DD3 9C0E C40A 488E C654 FB33 2AD5 9860
http://member.acm.org/~tomaszrybak
signature.asc
Description: This is a digitally signed message part
_______________________________________________ PyCUDA mailing list PyCUDA@tiker.net http://lists.tiker.net/listinfo/pycuda
