作者tmbyksdG (雨神妹妹的男朋友)
看板C_and_CPP
標題[問題] CUDA新手上路 平行度及硬體相關問題
時間Sun Nov 27 18:04:32 2016
開發平台(Platform): (Ex: Win10, Linux, ...)
Linux上安裝CUDA環境 （CUDA版本為8.0 運算能力為3.7）(Tesla K80)
編譯器(Ex: GCC, clang, VC++...)+目標環境(跟開發平台不同的話需列出)
NVCC
額外使用到的函數庫(Library Used): (Ex: OpenGL, ...)
問題(Question)：
硬體方面：
1.我執行deviceQuery偵測到2個device(device0:Tesla K80, device1:K80),估狗發現K80
是由兩個GK210核心所組成，那我偵測到的device是指有兩個K80（4個GK210）的意思嗎？
還是偵測到的兩個device其實就是GK210?
2.13個SMX,總共有2496個cores,所以我一次可以同時做運算的數量是否為2496個threads?
軟體方面：
我寫了一個64 * 64的矩陣乘法，我想測試不同的block & thread數量去做運算，哪個執
行時間會比較短，我試了兩種block的配置（thread數量剛好為4096，一個thread執行一
個輸出矩陣的一個element）。
（1）dim3 dimBlock(32, 32);
dim3 dimGrid(2, 2);
這個配置是以下附的程式碼配置，執行結果是正常的。
（2）dim3 dimBlock(4, 128);
dim3 dimGrid(1, 8);
換成這樣配置編譯之後，結果卻只有大約三分之一是正常的值，其餘卻都是0
，我的threadIdx.x & threadIdx.y都是由0-1023，請問這是發生了什麼錯誤嗎？
另外，我想知道執行運算時thread做了什麼事情，每個thread裡面裝的是什麼東西呢，是
我輸入的資料嗎？還是加跟乘這兩個指令呢？
餵入的資料(Input)：
預期的正確結果(Expected Output)：
錯誤結果(Wrong Output)：
程式碼(Code)：(請善用置底文網頁, 記得排版)
#include <stdio.h>
#include <stdlib.h>
#include <time.h>
#include <cuda_runtime.h>
#include "cuda_error.h"
#define N 64
__global__ void MatrixMulKernel(float* C, float* A, float* B, int n)
{
float Pvalue = 0;
int tx = threadIdx.x + blockIdx.x * blockDim.x;
int ty = threadIdx.y + blockIdx.y * blockDim.y;
for(int k = 0; k < n; k++)
{
float Aelement = A[ty * n + k];
float Belement = B[k * n + tx];
Pvalue += Aelement * Belement;
}
C[ty * n + tx] = Pvalue;
__syncthreads();
}
int main()
{
float A[N][N];
float B[N][N];
float C[N][N];
size_t size = N * N * sizeof(float);
int i, j;
float *gA, *gB, *gC;
for(i = 0; i < N; i++)
{
for(j = 0; j < N; j++)
{
A[i][j] = (float) (1 + ( rand() % 9 ));
B[i][j] = (float) (1 + ( rand() % 9 ));
}
}
/*allocate memory*/
cudaMalloc(&gA, size);
cudaMalloc(&gB, size);
cudaMalloc(&gC, size);
puts(cudaGetErrorString(cudaGetLastError()));
cudaMemcpy(gA, A, size, cudaMemcpyHostToDevice);
cudaMemcpy(gB, B, size, cudaMemcpyHostToDevice);
puts(cudaGetErrorString(cudaGetLastError()));
float etime;
cudaEvent_t start, stop;
cudaEventCreate(&start, cudaEventDefault);
cudaEventCreate(&stop, cudaEventDefault);
dim3 dimBlock( 32, 32 );
dim3 dimGrid( 2, 2 );
cudaEventRecord(start);
/*call GPU kernel*/
MatrixMulKernel<<<dimGrid,dimBlock>>>(gC, gA, gB, N);
cudaEventRecord(stop);
cudaEventSynchronize(stop);
cudaEventElapsedTime(&etime, start, stop);
printf("%f ms\n", etime);
puts(cudaGetErrorString(cudaGetLastError()));
cudaMemcpy(C, gC, size, cudaMemcpyDeviceToHost);
puts(cudaGetErrorString(cudaGetLastError()));
cudaFree(gA);
cudaFree(gB);
cudaFree(gC);
return 0;
}
補充說明(Supplement)：
請版上的各位先進指導一下我，謝謝。另外，手機排版請見諒。
--
※ 發信站: 批踢踢實業坊(ptt.cc), 來自: 59.115.193.128
※ 文章網址: https://webptt.com/m.aspx?n=bbs/C_and_CPP/M.1480241075.A.D3D.html
1F：→ opl164: 第二個狀況的tx,ty 你可以印出來看看跟的一個會不一樣吧 11/28 12:03
2F：推 a1u1usul3: 硬體1應該是兩個GK210的意思 11/28 14:31
3F：→ a1u1usul3: 硬體2可以說是對的，精確來說是78個wrap吧 11/28 14:32
4F：→ a1u1usul3: 軟體問題可以用cuda-memcheck跑看看，是不是access到奇 11/28 14:34
5F：→ a1u1usul3: 怪的地方去了 11/28 14:34
6F：→ arXiv: 第二點錯了程式跟著錯,請先讀文件cuda-c-programming-guide 11/29 00:09