[问题] 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