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

作者: tmbyksdG (雨神妹妹的男朋友)   2016-11-27 18:04:32
開發平台(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):
請版上的各位先進指導一下我,謝謝。另外,手機排版請見諒。
作者: opl164 (opl)   2016-11-28 12:03:00
第二個狀況的tx,ty 你可以印出來看看跟的一個會不一樣吧
作者: a1u1usul3 (Q-Max)   2016-11-28 14:31:00
硬體1應該是兩個GK210的意思硬體2可以說是對的,精確來說是78個wrap吧軟體問題可以用cuda-memcheck跑看看,是不是access到奇怪的地方去了
作者: arXiv (arXiv)   2016-11-29 00:09:00
第二點錯了程式跟著錯,請先讀文件cuda-c-programming-guide

Links booklink

Contact Us: admin [ a t ] ucptt.com