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;
}
}
}
parallel-processing cuda gpu gpgpu gpu-programming
add a comment |
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;
}
}
}
parallel-processing cuda gpu gpgpu gpu-programming
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 unrollnby 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
add a comment |
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;
}
}
}
parallel-processing cuda gpu gpgpu gpu-programming
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
parallel-processing cuda gpu gpgpu gpu-programming
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 unrollnby 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
add a comment |
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 unrollnby 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
add a comment |
active
oldest
votes
active
oldest
votes
active
oldest
votes
active
oldest
votes
active
oldest
votes
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.
Sign up or log in
StackExchange.ready(function () {
StackExchange.helpers.onClickDraftSave('#login-link');
});
Sign up using Google
Sign up using Facebook
Sign up using Email and Password
Post as a guest
Required, but never shown
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
Sign up or log in
StackExchange.ready(function () {
StackExchange.helpers.onClickDraftSave('#login-link');
});
Sign up using Google
Sign up using Facebook
Sign up using Email and Password
Post as a guest
Required, but never shown
Sign up or log in
StackExchange.ready(function () {
StackExchange.helpers.onClickDraftSave('#login-link');
});
Sign up using Google
Sign up using Facebook
Sign up using Email and Password
Post as a guest
Required, but never shown
Sign up or log in
StackExchange.ready(function () {
StackExchange.helpers.onClickDraftSave('#login-link');
});
Sign up using Google
Sign up using Facebook
Sign up using Email and Password
Sign up using Google
Sign up using Facebook
Sign up using Email and Password
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
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
nby 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