[問題] CUDA新手上路 平行度及硬體相關問題

[tmbyksdG (雨神妹妹的男朋友)]

開發平台(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://www.ptt.cc/bbs/C_and_CPP/M.1480241075.A.D3D.html
→ opl164: 第二個狀況的tx,ty 你可以印出來看看跟的一個會不一樣吧
推 a1u1usul3: 硬體1應該是兩個GK210的意思
→ a1u1usul3: 硬體2可以說是對的，精確來說是78個wrap吧
→ a1u1usul3: 軟體問題可以用cuda-memcheck跑看看，是不是access到奇
→ a1u1usul3: 怪的地方去了
→ arXiv: 第二點錯了程式跟著錯,請先讀文件cuda-c-programming-guide