쿠다 프로그래밍에서 float와 double 유형의 차이점은 무엇입니까?

Question

전 cuda 프로그래밍에 익숙합니다. 내 프로그램 (공유 메모리를 사용하는 행렬 곱셈)에서 block_size = 20을 정의했고 행렬이 1200 * 1200 일 때 프로그램은 double 요소와 작동하지만 float 요소에서는 작동하지 않습니다 (요소가 부동 상태 인 경우 840 * 840 행렬에서 작동 함). 내 질문은 그것이 일어나는 이유입니다, 우리는 float 타입이 double보다 작다는 것을 알고 있습니다 만?

// Matrices are stored in row-major order: 
// M(row, col) = *(M.elements + row * M.stride + col) 
#include <stdio.h> 
#define BLOCK_SIZE 20 
typedef struct { 
int width; 
int height; 
int stride; 
float* elements; 
} Matrix; 
// Get a matrix element 
__device__ float GetElement(const Matrix A, int row, int col) 
{ 
return A.elements[row * A.stride + col]; 
} 
// Set a matrix element 
__device__ void SetElement(Matrix A, int row, int col, 
float value) 
{ 
A.elements[row * A.stride + col] = value; 
} 
// Get the BLOCK_SIZExBLOCK_SIZE sub-matrix Asub of A that is 
// located col sub-matrices to the right and row sub-matrices down 
// from the upper-left corner of A 
__device__ Matrix GetSubMatrix(Matrix A, int row, int col) 
{ 
Matrix Asub; 

Asub.width = BLOCK_SIZE; 
Asub.height = BLOCK_SIZE; 
Asub.stride = A.stride; 
Asub.elements = &A.elements[A.stride * BLOCK_SIZE * row+ BLOCK_SIZE * col]; 
return Asub; 
} 
// Thread block size 
// Forward declaration of the matrix multiplication kernel 
__global__ void MatMulKernel(const Matrix, const Matrix, Matrix); 
// Matrix multiplication - Host code 
// Matrix dimensions are assumed to be multiples of BLOCK_SIZE 
void MatMul(const Matrix A, const Matrix B, Matrix C) 
{ 

// Load A and B to device memory 
Matrix d_A; 
d_A.width = d_A.stride = A.width; d_A.height = A.height; 
siz e_t size = A.width * A.height * sizeof(float); 
cudaMalloc((void **)&d_A.elements, size); 
cudaMemcpy(d_A.elements, A.elements, size, 
cudaMemcpyHostToDevice); 
Matrix d_B; 
d_B.width = d_B.stride = B.width; d_B.height = B.height; 
size = B.width * B.height * sizeof(float); 
cudaMalloc((void **)&d_B.elements, size); 
cudaMemcpy(d_B.elements, B.elements, size, 
cudaMemcpyHostToDevice); 
// Allocate C in device memory 
Matrix d_C; 
d_C.width = d_C.stride = C.width; d_C.height = C.height; 
size = C.width * C.height * sizeof(float); 
cudaMalloc((void **)&d_C.elements, size); 
// Invoke kernel 
dim3 dimBlock(BLOCK_SIZE,BLOCK_SIZE); 
//dim3 dimBlock(C.height, C.width); 
//dim3 dimGrid(B.width/dimBlock.x, A.height/dimBlock.y); 
dim3 dimGrid((B.width+dimBlock.x-1)/dimBlock.x, (A.height+dimBlock.y-1) /dimBlock.y); 
MatMulKernel<<<dimGrid, dimBlock>>>(d_A, d_B, d_C); 
// Read C from device memory 
cudaMemcpy(C.elements, d_C.elements, size, 
cudaMemcpyDeviceToHost); 
// Free device memory 
cudaFree(d_A.elements); 
cudaFree(d_B.elements); 
cudaFree(d_C.elements); 
} 
// Matrix multiplication kernel called by MatMul() 
__global__ void MatMulKernel(Matrix A, Matrix B, Matrix C) 
{ 
// Block row and column 
int blockRow = blockIdx.y; 
int blockCol = blockIdx.x; 
// Each thread block computes one sub-matrix Csub of C 
Matrix Csub = GetSubMatrix(C, blockRow, blockCol); 
// Each thread computes one element of Csub 
// by accumulating results into Cvalue 
float Cvalue = 0; 
// Thread row and column within Csub 
int row = threadIdx.y; 
int col = threadIdx.x; 
// Loop over all the sub-matrices of A and B that are 
// required to compute Csub 
// Multiply each pair of sub-matrices together 
// and accumulate the results 
for (int m = 0; m < (A.width/BLOCK_SIZE); ++m) { 
// Get sub-matrix Asub of A 
Matrix Asub = GetSubMatrix(A, blockRow, m); 
// Get sub-matrix Bsub of B 
Matrix Bsub = GetSubMatrix(B, m, blockCol); 
// Shared memory used to store Asub and Bsub respectively 
__shared__ float As[BLOCK_SIZE][BLOCK_SIZE]; 
__shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE]; 
// Load Asub and Bsub from device memory to shared memory 
// Each thread loads one element of each sub-matrix 
As[row][col] = GetElement(Asub, row, col); 
Bs[row][col] = GetElement(Bsub, row, col); 
// Synchronize to make sure the sub-matrices are loaded 
// before starting the computation 
__syncthreads(); 
// Multiply Asub and Bsub together 
for (int e = 0; e < BLOCK_SIZE; ++e) 
Cvalue += As[row][e] * Bs[e][col]; 
// Synchronize to make sure that the preceding 
// computation is done before loading two new 
// sub-matrices of A and B in the next iteration 
__syncthreads(); 
} 
// Write Csub to device memory 
// Each thread writes one element 
SetElement(Csub, row, col, Cvalue); 
} 
////////////////////////////////////////////////////////// 
/// print_matrix function /////////////////////////// 
//////////////////////////////////////////////////////// 
void print_matrix(float *c,int row,int col){ 
for (int i = 0; i < row; ++i){ 
for (int j = 0; j < col; ++j) 
printf("%f ",c[col*i +j]); 
printf("\n\n"); 
} 
} 
////////////////////////////////////////////////////////// 
/// random_init function /////////////////////////// 
//////////////////////////////////////////////////////// 
void random_init(float *a,int size){ 
for(int i=0;i<size;i++) 
a[i]=rand()%10; 
} 
//////////////////////////////////////////////////////// 
int main(void){ 

//////////////////////////////////////////////////////\|/ 
cudaEvent_t start,stop; 
///////////////////////////////////////////////////////|\ 

Matrix A,B,C; 
A.width=1200; 
A.height=1200;///// 
B.width=1200;///// 
B.height=1200; 
C.width=B.width; 
C.height=A.height; 

size_t size = A.width * A.height * sizeof(float); 
A.elements = (float *)malloc(size); 
//random_init(A.elements,A.width * A.height); 
size = B.width * B.height * sizeof(float); 
B.elements= (float *)malloc(size); 
//random_init(B.elements,B.width * B.height); 
size = C.width * C.height * sizeof(float); 
C.elements= (float *)malloc(size); 
for(int i=0;i<A.width*A.height;i++) 
A.elements[i]=1; 
for(int i=0;i<B.width*B.height;i++) 
B.elements[i]=1; 
printf("matrix A(%d,%d) & matrix B(%d,%d) & matrix C(%d,%d)\n",A.width,A.height,B.width, 
B.height,C.width,C.height); 
//////////////////////////////////////////////////////\|/ 
cudaEventCreate(&start); 
cudaEventCreate(&stop); 
cudaEventRecord(start,0); 
///////////////////////////////////////////////////////|\ 

MatMul(A,B,C); 
//////////////////////////////////////////////////////\|/ 
cudaEventRecord(stop,0); 
cudaEventSynchronize(stop); 
float elapsedTime; 
cudaEventElapsedTime(&elapsedTime,start,stop); 
printf("Time to genreat : %3.5f ms\n",elapsedTime); 
///////////////////////////////////////////////////////|\ 
printf("\nC\n"); 
//print_matrix(C.elements,C.height,C.width); 


printf("C[%d]=%f\n",0,C.elements[0]); 
printf("C[%d]=%f\n",C.width -1,C.elements[C.width-1]); 
printf("C[%d]=%f\n",(C.width * C.height)-1,C.elements[(C.width * C.height)-1]); 

getchar(); 
return(0); 
} 

Answer 1

다음과 같은 메시지가 :

""디스플레이 드라이버가 응답을 중지하고 복구되었습니다 "

당신이 windows TDR event으로 실행 한 표시입니다."

Windows에서 실행하기에는 너무 오래 걸리는 커널은 윈도우 디스플레이 워치 독 타이머가 디스플레이 장치를 재설정하게하여 CUDA 코드 실행을 종료시킵니다. 실행하는데 약 2 초 이상을 필요로하는 커널이 실행될 수 있습니다.

"windows TDR"을 검색하면 다른 설명과이 문제를 해결할 수있는 방법이 있습니다. 또한 변경 한 코드가 실행되는 데 왜 오래 걸리는지 조사 할 수 있습니다.