StackOverflow2013

Note that there are some explanatory texts on larger screens.

plurals

PO
primarykey
Id
7191690
data
AcceptedAnswerId
0
AnswerCount
0
ClosedDate
CommentCount
0
CommunityOwnedDate
CreationDate
2011-08-25T13:59:42.540
FavoriteCount
0
LastActivityDate
2011-08-25T13:59:42.540
LastEditDate
LastEditorUserId
0
OwnerUserId
635654
ParentId
7164210
PostTypeId
2
Score
0
ViewCount
0
LastEditorDisplayName
text
Body
While there is no official answer to this, you can measure through atomic operations when your blocks begin your work and when they end. Try playing with the following code: <pre><code>#include <stdio.h> const int maxBlocks=60; //Number of blocks of size 512 threads on current device required to achieve full occupancy __global__ void emptyKernel() {} __global__ void myKernel(int *control, int *output) { if (threadIdx.x==1) { //register that we enter int enter=atomicAdd(control,1); output[blockIdx.x]=enter; //some intensive and long task int &var=output[blockIdx.x+gridDim.x]; //var references global memory var=1; for (int i=0; i<12345678; ++i) { var+=1+tanhf(var); } //register that we quit var=atomicAdd(control,1); } } int main() { int *gpuControl; cudaMalloc((void**)&gpuControl, sizeof(int)); int cpuControl=0; cudaMemcpy(gpuControl,&cpuControl,sizeof(int),cudaMemcpyHostToDevice); int *gpuOutput; cudaMalloc((void**)&gpuOutput, sizeof(int)*maxBlocks*2); int cpuOutput[maxBlocks*2]; for (int i=0; i<maxBlocks*2; ++i) //clear the host array just to be on the safe side cpuOutput[i]=-1; // play with these values const int thr=479; const int p=13; const int q=maxBlocks; //I found that this may actually affect the scheduler! Try with and without this call. emptyKernel<<<p,thr>>>(); cudaEvent_t timerStart; cudaEvent_t timerStop; cudaEventCreate(&timerStart); cudaEventCreate(&timerStop); cudaThreadSynchronize(); cudaEventRecord(timerStart,0); myKernel<<<q,512>>>(gpuControl, gpuOutput); cudaEventRecord(timerStop,0); cudaEventSynchronize(timerStop); cudaMemcpy(cpuOutput,gpuOutput,sizeof(int)*maxBlocks*2,cudaMemcpyDeviceToHost); cudaThreadSynchronize(); float thisTime; cudaEventElapsedTime(&thisTime,timerStart,timerStop); cudaEventDestroy(timerStart); cudaEventDestroy(timerStop); printf("Elapsed time: %f\n",thisTime); for (int i=0; i<q; ++i) printf("%d: %d-%d\n",i,cpuOutput[i],cpuOutput[i+q]); } </code></pre> What you get in the output is the block ID, followed by the enter "time" and exit "time". This way you can learn in which order those events occured.
Tags
Title
singulars
PostAcceptedAnswerId
1. This table or related slice is empty.
PostParentId
1. POA question about the details about the distribution from blocks to SMs in CUDA
 singulars
 PostTypePostTypeId
 PTQuestion
PostTypePostTypeId
1. PTAnswer
UserLastEditorUserId
1. This table or related slice is empty.
UserOwnerUserId
1. USCygnusX1
plurals
PostLinksPostIdRelatedPostId
1. This table or related slice is empty.
PostLinksRelatedPostIdPostId
1. This table or related slice is empty.
PostsAcceptedAnswerId
1. This table or related slice is empty.
PostsParentIdCreationDate
1. This table or related slice is empty.
VotesPostIdCreationDate
1. This table or related slice is empty.
CommentsPostId
1. This table or related slice is empty.

Querying!

Guidance

A row detail

Detail views are divided into sections. All the information in the data section comes from columns in the selected row. The other sections display data from other, related rows.

Related data can be related in a to-one or a to-many fashion. Captions of data related in a to-many fashion link to a list view showing a filtered view of the table.

Try moving around until you find a non-empty to-many entry and click on the label to get to one. You can move back to the root by clicking on the database name in the header.