CUDA same execution time for larger thread tasks











up vote
-2
down vote

favorite












I created a program in CUDA to multiply two matrices together. I used multiple blocks and also increased the amount of elements each thread computes. I measured the execution time for multiple block sizes and multiple workloads for each thread, but I'm a little confused by the results. Changing the block size changes the execution time, but changing how many elements each thread should compute doesn't change the execution time at all. I'm not sure if there's something wrong with my code or if I am just misunderstanding the problem. I would think that as the number of elements each thread must compute increases, the execution time should increase as well. I am getting correct output however so I don't understand what could be wrong.



In my code, TILE_WIDTH (B) represents the block size and W_WIDTH (W) represents how many elements each thread should compute. So, a BxB block should compute BWxBW elements with each thread computing WxW elements.



Here's the code to invoke the kernel. As W increases, the number of blocks decreases. Width is the size of one side of a matrix (4096).



    dim3 dimGrid(Width/(TILE_WIDTH*W_WIDTH), Width/(TILE_WIDTH*W_WIDTH));
dim3 dimBlock(TILE_WIDTH, TILE_WIDTH);
MatrixMulKernel_sharedmemory<<<dimGrid, dimBlock>>>(ad, bd, cd, Width);


Here's the kernel:



__global__ void MatrixMulKernel_sharedmemory(float* ad, float* bd, float* cd, int Width) {
__shared__ float ads[TILE_WIDTH][TILE_WIDTH];
__shared__ float bds[TILE_WIDTH][TILE_WIDTH];

int Bx = blockIdx.x;
int By = blockIdx.y;

int Tx = threadIdx.x;
int Ty = threadIdx.y;

int rowd = By * TILE_WIDTH * W_WIDTH + Ty;
int columnd = Bx * TILE_WIDTH * W_WIDTH+ Tx;
int rowds;
int columnds;

for (int m = 0; m < W_WIDTH; m++) {
for (int n = 0; n < W_WIDTH; n++) {
rowds = rowd + m*TILE_WIDTH;
columnds = columnd + n*TILE_WIDTH;
float Cvalue = 0;

for(int i = 0; i < Width /( TILE_WIDTH); i++) {
ads[Ty][Tx] = ad[rowds * Width + i * TILE_WIDTH + Tx];
bds[Ty][Tx] = bd[(i*TILE_WIDTH + Ty)*Width + columnds];
__syncthreads();


for (int k = 0; k < TILE_WIDTH; k++) {
Cvalue += ads[Ty][k] * bds[k][Tx];
__syncthreads();
}
}

cd[rowds * Width + columnds] = Cvalue;

}
}
}









share|improve this question






















  • Do you need syncthreads in the last for loop? What happens if you carry that syncthreads to just after the for-loop? Perhaps too many synching in that last for loop causes the most latency so that computing more elements per kernel increases that latency? Isn't doing more work per kernel is to harness ILP with broken chain of dependencies? Maybe you should try unroll n by 2 so that each syncthread synchronizes 2 items at once, increasing ILP between two computes? Perhaps you had good TLP but increasing number of items per kernel increased register pressure and decreased TLP?
    – huseyin tugrul buyukisik
    Nov 11 at 19:21

















up vote
-2
down vote

favorite












I created a program in CUDA to multiply two matrices together. I used multiple blocks and also increased the amount of elements each thread computes. I measured the execution time for multiple block sizes and multiple workloads for each thread, but I'm a little confused by the results. Changing the block size changes the execution time, but changing how many elements each thread should compute doesn't change the execution time at all. I'm not sure if there's something wrong with my code or if I am just misunderstanding the problem. I would think that as the number of elements each thread must compute increases, the execution time should increase as well. I am getting correct output however so I don't understand what could be wrong.



In my code, TILE_WIDTH (B) represents the block size and W_WIDTH (W) represents how many elements each thread should compute. So, a BxB block should compute BWxBW elements with each thread computing WxW elements.



Here's the code to invoke the kernel. As W increases, the number of blocks decreases. Width is the size of one side of a matrix (4096).



    dim3 dimGrid(Width/(TILE_WIDTH*W_WIDTH), Width/(TILE_WIDTH*W_WIDTH));
dim3 dimBlock(TILE_WIDTH, TILE_WIDTH);
MatrixMulKernel_sharedmemory<<<dimGrid, dimBlock>>>(ad, bd, cd, Width);


Here's the kernel:



__global__ void MatrixMulKernel_sharedmemory(float* ad, float* bd, float* cd, int Width) {
__shared__ float ads[TILE_WIDTH][TILE_WIDTH];
__shared__ float bds[TILE_WIDTH][TILE_WIDTH];

int Bx = blockIdx.x;
int By = blockIdx.y;

int Tx = threadIdx.x;
int Ty = threadIdx.y;

int rowd = By * TILE_WIDTH * W_WIDTH + Ty;
int columnd = Bx * TILE_WIDTH * W_WIDTH+ Tx;
int rowds;
int columnds;

for (int m = 0; m < W_WIDTH; m++) {
for (int n = 0; n < W_WIDTH; n++) {
rowds = rowd + m*TILE_WIDTH;
columnds = columnd + n*TILE_WIDTH;
float Cvalue = 0;

for(int i = 0; i < Width /( TILE_WIDTH); i++) {
ads[Ty][Tx] = ad[rowds * Width + i * TILE_WIDTH + Tx];
bds[Ty][Tx] = bd[(i*TILE_WIDTH + Ty)*Width + columnds];
__syncthreads();


for (int k = 0; k < TILE_WIDTH; k++) {
Cvalue += ads[Ty][k] * bds[k][Tx];
__syncthreads();
}
}

cd[rowds * Width + columnds] = Cvalue;

}
}
}









share|improve this question






















  • Do you need syncthreads in the last for loop? What happens if you carry that syncthreads to just after the for-loop? Perhaps too many synching in that last for loop causes the most latency so that computing more elements per kernel increases that latency? Isn't doing more work per kernel is to harness ILP with broken chain of dependencies? Maybe you should try unroll n by 2 so that each syncthread synchronizes 2 items at once, increasing ILP between two computes? Perhaps you had good TLP but increasing number of items per kernel increased register pressure and decreased TLP?
    – huseyin tugrul buyukisik
    Nov 11 at 19:21















up vote
-2
down vote

favorite









up vote
-2
down vote

favorite











I created a program in CUDA to multiply two matrices together. I used multiple blocks and also increased the amount of elements each thread computes. I measured the execution time for multiple block sizes and multiple workloads for each thread, but I'm a little confused by the results. Changing the block size changes the execution time, but changing how many elements each thread should compute doesn't change the execution time at all. I'm not sure if there's something wrong with my code or if I am just misunderstanding the problem. I would think that as the number of elements each thread must compute increases, the execution time should increase as well. I am getting correct output however so I don't understand what could be wrong.



In my code, TILE_WIDTH (B) represents the block size and W_WIDTH (W) represents how many elements each thread should compute. So, a BxB block should compute BWxBW elements with each thread computing WxW elements.



Here's the code to invoke the kernel. As W increases, the number of blocks decreases. Width is the size of one side of a matrix (4096).



    dim3 dimGrid(Width/(TILE_WIDTH*W_WIDTH), Width/(TILE_WIDTH*W_WIDTH));
dim3 dimBlock(TILE_WIDTH, TILE_WIDTH);
MatrixMulKernel_sharedmemory<<<dimGrid, dimBlock>>>(ad, bd, cd, Width);


Here's the kernel:



__global__ void MatrixMulKernel_sharedmemory(float* ad, float* bd, float* cd, int Width) {
__shared__ float ads[TILE_WIDTH][TILE_WIDTH];
__shared__ float bds[TILE_WIDTH][TILE_WIDTH];

int Bx = blockIdx.x;
int By = blockIdx.y;

int Tx = threadIdx.x;
int Ty = threadIdx.y;

int rowd = By * TILE_WIDTH * W_WIDTH + Ty;
int columnd = Bx * TILE_WIDTH * W_WIDTH+ Tx;
int rowds;
int columnds;

for (int m = 0; m < W_WIDTH; m++) {
for (int n = 0; n < W_WIDTH; n++) {
rowds = rowd + m*TILE_WIDTH;
columnds = columnd + n*TILE_WIDTH;
float Cvalue = 0;

for(int i = 0; i < Width /( TILE_WIDTH); i++) {
ads[Ty][Tx] = ad[rowds * Width + i * TILE_WIDTH + Tx];
bds[Ty][Tx] = bd[(i*TILE_WIDTH + Ty)*Width + columnds];
__syncthreads();


for (int k = 0; k < TILE_WIDTH; k++) {
Cvalue += ads[Ty][k] * bds[k][Tx];
__syncthreads();
}
}

cd[rowds * Width + columnds] = Cvalue;

}
}
}









share|improve this question













I created a program in CUDA to multiply two matrices together. I used multiple blocks and also increased the amount of elements each thread computes. I measured the execution time for multiple block sizes and multiple workloads for each thread, but I'm a little confused by the results. Changing the block size changes the execution time, but changing how many elements each thread should compute doesn't change the execution time at all. I'm not sure if there's something wrong with my code or if I am just misunderstanding the problem. I would think that as the number of elements each thread must compute increases, the execution time should increase as well. I am getting correct output however so I don't understand what could be wrong.



In my code, TILE_WIDTH (B) represents the block size and W_WIDTH (W) represents how many elements each thread should compute. So, a BxB block should compute BWxBW elements with each thread computing WxW elements.



Here's the code to invoke the kernel. As W increases, the number of blocks decreases. Width is the size of one side of a matrix (4096).



    dim3 dimGrid(Width/(TILE_WIDTH*W_WIDTH), Width/(TILE_WIDTH*W_WIDTH));
dim3 dimBlock(TILE_WIDTH, TILE_WIDTH);
MatrixMulKernel_sharedmemory<<<dimGrid, dimBlock>>>(ad, bd, cd, Width);


Here's the kernel:



__global__ void MatrixMulKernel_sharedmemory(float* ad, float* bd, float* cd, int Width) {
__shared__ float ads[TILE_WIDTH][TILE_WIDTH];
__shared__ float bds[TILE_WIDTH][TILE_WIDTH];

int Bx = blockIdx.x;
int By = blockIdx.y;

int Tx = threadIdx.x;
int Ty = threadIdx.y;

int rowd = By * TILE_WIDTH * W_WIDTH + Ty;
int columnd = Bx * TILE_WIDTH * W_WIDTH+ Tx;
int rowds;
int columnds;

for (int m = 0; m < W_WIDTH; m++) {
for (int n = 0; n < W_WIDTH; n++) {
rowds = rowd + m*TILE_WIDTH;
columnds = columnd + n*TILE_WIDTH;
float Cvalue = 0;

for(int i = 0; i < Width /( TILE_WIDTH); i++) {
ads[Ty][Tx] = ad[rowds * Width + i * TILE_WIDTH + Tx];
bds[Ty][Tx] = bd[(i*TILE_WIDTH + Ty)*Width + columnds];
__syncthreads();


for (int k = 0; k < TILE_WIDTH; k++) {
Cvalue += ads[Ty][k] * bds[k][Tx];
__syncthreads();
}
}

cd[rowds * Width + columnds] = Cvalue;

}
}
}






parallel-processing cuda gpu gpgpu gpu-programming






share|improve this question













share|improve this question











share|improve this question




share|improve this question










asked Nov 9 at 7:50









Tyler Robins

121




121












  • Do you need syncthreads in the last for loop? What happens if you carry that syncthreads to just after the for-loop? Perhaps too many synching in that last for loop causes the most latency so that computing more elements per kernel increases that latency? Isn't doing more work per kernel is to harness ILP with broken chain of dependencies? Maybe you should try unroll n by 2 so that each syncthread synchronizes 2 items at once, increasing ILP between two computes? Perhaps you had good TLP but increasing number of items per kernel increased register pressure and decreased TLP?
    – huseyin tugrul buyukisik
    Nov 11 at 19:21




















  • Do you need syncthreads in the last for loop? What happens if you carry that syncthreads to just after the for-loop? Perhaps too many synching in that last for loop causes the most latency so that computing more elements per kernel increases that latency? Isn't doing more work per kernel is to harness ILP with broken chain of dependencies? Maybe you should try unroll n by 2 so that each syncthread synchronizes 2 items at once, increasing ILP between two computes? Perhaps you had good TLP but increasing number of items per kernel increased register pressure and decreased TLP?
    – huseyin tugrul buyukisik
    Nov 11 at 19:21


















Do you need syncthreads in the last for loop? What happens if you carry that syncthreads to just after the for-loop? Perhaps too many synching in that last for loop causes the most latency so that computing more elements per kernel increases that latency? Isn't doing more work per kernel is to harness ILP with broken chain of dependencies? Maybe you should try unroll n by 2 so that each syncthread synchronizes 2 items at once, increasing ILP between two computes? Perhaps you had good TLP but increasing number of items per kernel increased register pressure and decreased TLP?
– huseyin tugrul buyukisik
Nov 11 at 19:21






Do you need syncthreads in the last for loop? What happens if you carry that syncthreads to just after the for-loop? Perhaps too many synching in that last for loop causes the most latency so that computing more elements per kernel increases that latency? Isn't doing more work per kernel is to harness ILP with broken chain of dependencies? Maybe you should try unroll n by 2 so that each syncthread synchronizes 2 items at once, increasing ILP between two computes? Perhaps you had good TLP but increasing number of items per kernel increased register pressure and decreased TLP?
– huseyin tugrul buyukisik
Nov 11 at 19:21



















active

oldest

votes











Your Answer






StackExchange.ifUsing("editor", function () {
StackExchange.using("externalEditor", function () {
StackExchange.using("snippets", function () {
StackExchange.snippets.init();
});
});
}, "code-snippets");

StackExchange.ready(function() {
var channelOptions = {
tags: "".split(" "),
id: "1"
};
initTagRenderer("".split(" "), "".split(" "), channelOptions);

StackExchange.using("externalEditor", function() {
// Have to fire editor after snippets, if snippets enabled
if (StackExchange.settings.snippets.snippetsEnabled) {
StackExchange.using("snippets", function() {
createEditor();
});
}
else {
createEditor();
}
});

function createEditor() {
StackExchange.prepareEditor({
heartbeatType: 'answer',
convertImagesToLinks: true,
noModals: true,
showLowRepImageUploadWarning: true,
reputationToPostImages: 10,
bindNavPrevention: true,
postfix: "",
imageUploader: {
brandingHtml: "Powered by u003ca class="icon-imgur-white" href="https://imgur.com/"u003eu003c/au003e",
contentPolicyHtml: "User contributions licensed under u003ca href="https://creativecommons.org/licenses/by-sa/3.0/"u003ecc by-sa 3.0 with attribution requiredu003c/au003e u003ca href="https://stackoverflow.com/legal/content-policy"u003e(content policy)u003c/au003e",
allowUrls: true
},
onDemand: true,
discardSelector: ".discard-answer"
,immediatelyShowMarkdownHelp:true
});


}
});














draft saved

draft discarded


















StackExchange.ready(
function () {
StackExchange.openid.initPostLogin('.new-post-login', 'https%3a%2f%2fstackoverflow.com%2fquestions%2f53221700%2fcuda-same-execution-time-for-larger-thread-tasks%23new-answer', 'question_page');
}
);

Post as a guest















Required, but never shown






























active

oldest

votes













active

oldest

votes









active

oldest

votes






active

oldest

votes
















draft saved

draft discarded




















































Thanks for contributing an answer to Stack Overflow!


  • Please be sure to answer the question. Provide details and share your research!

But avoid



  • Asking for help, clarification, or responding to other answers.

  • Making statements based on opinion; back them up with references or personal experience.


To learn more, see our tips on writing great answers.





Some of your past answers have not been well-received, and you're in danger of being blocked from answering.


Please pay close attention to the following guidance:


  • Please be sure to answer the question. Provide details and share your research!

But avoid



  • Asking for help, clarification, or responding to other answers.

  • Making statements based on opinion; back them up with references or personal experience.


To learn more, see our tips on writing great answers.




draft saved


draft discarded














StackExchange.ready(
function () {
StackExchange.openid.initPostLogin('.new-post-login', 'https%3a%2f%2fstackoverflow.com%2fquestions%2f53221700%2fcuda-same-execution-time-for-larger-thread-tasks%23new-answer', 'question_page');
}
);

Post as a guest















Required, but never shown





















































Required, but never shown














Required, but never shown












Required, but never shown







Required, but never shown

































Required, but never shown














Required, but never shown












Required, but never shown







Required, but never shown







這個網誌中的熱門文章

Academy of Television Arts & Sciences

L'Équipe

1995 France bombings