Hello all,
I am using a trial version of pgi v14.4. I have a simple matrix multiply code that I am attempting to parallelize using openacc. I have cuda 5.5 installed on a linux machine.
Here is the sample code:
/* matrix-acc-check.c */
#define SIZE 2000
float a[SIZE][SIZE];
float b[SIZE][SIZE];
float c[SIZE][SIZE];
float seq[SIZE][SIZE];
int main()
{
int i,j,k;
// Initialize matrices.
for (i = 0; i < SIZE; ++i) {
for (j = 0; j < SIZE; ++j) {
a[i][j] = (float)i + j;
b[i][j] = (float)i - j;
c[i][j] = 0.0f;
}
}
// Compute matrix multiplication.
#pragma acc kernels copyin(a,b) copy(c)
for (i = 0; i < SIZE; ++i) {
for (j = 0; j < SIZE; ++j) {
for (k = 0; k < SIZE; ++k) {
c[i][j] += a[i][k] * b[k][j];
}
}
}
// ****************
// double-check the OpenACC result sequentially on the host
// ****************
// Initialize the seq matrix
for(i = 0; i < SIZE; ++i)
for(j = 0; j < SIZE; ++j)
seq[i][j] = 0.f;
// Perform the multiplication
for (i = 0; i < SIZE; ++i)
for (j = 0; j < SIZE; ++j)
for (k = 0; k < SIZE; ++k)
seq[i][j] += a[i][k] * b[k][j];
// check all the OpenACC matrices
for (i = 0; i < SIZE; ++i)
for (j = 0; j < SIZE; ++j)
if(c[i][j] != seq[i][j]) {
printf("Error %d %d\n", i,j);
exit(1);
}
printf("OpenACC matrix multiplication test was successful!\n");
return 0;
}
I compile it as follows:
pgcc -acc -fast -Minfo -ta=nvidia,cc1x matrix-mul.c
When the size of the matrix is set to 1000, the program runs fine. However, when the size of the matrix is set to 2000, I get the following error:
call to cuEventSynchronize returned error 702: Launch timeout
Does anyone provide an idea what is going wrong? I appreciate your help.
FYI, here is the output of devicequery:
./deviceQuery Starting…
CUDA Device Query (Runtime API) version (CUDART static linking)
Detected 1 CUDA Capable device(s)
Device 0: “Quadro FX 770M”
CUDA Driver Version / Runtime Version 5.5 / 5.0
CUDA Capability Major/Minor version number: 1.1
Total amount of global memory: 511 MBytes (536150016 bytes)
( 4) Multiprocessors x ( 8) CUDA Cores/MP: 32 CUDA Cores
GPU Clock rate: 1250 MHz (1.25 GHz)
Memory Clock rate: 800 Mhz
Memory Bus Width: 128-bit
Max Texture Dimension Size (x,y,z) 1D=(8192), 2D=(65536,32768), 3D=(2048,2048,2048)
Max Layered Texture Size (dim) x layers 1D=(8192) x 512, 2D=(8192,8192) x 512
Total amount of constant memory: 65536 bytes
Total amount of shared memory per block: 16384 bytes
Total number of registers available per block: 8192
Warp size: 32
Maximum number of threads per multiprocessor: 768
Maximum number of threads per block: 512
Maximum sizes of each dimension of a block: 512 x 512 x 64
Maximum sizes of each dimension of a grid: 65535 x 65535 x 1
Maximum memory pitch: 2147483647 bytes
Texture alignment: 256 bytes
Concurrent copy and kernel execution: Yes with 1 copy engine(s)
Run time limit on kernels: Yes
Integrated GPU sharing Host Memory: No
Support host page-locked memory mapping: Yes
Alignment requirement for Surfaces: Yes
Device has ECC support: Disabled
Device supports Unified Addressing (UVA): No
Device PCI Bus ID / PCI location ID: 1 / 0
Compute Mode:
< Default (multiple host threads can use ::cudaSetDevice() with device simultaneously) >deviceQuery, CUDA Driver = CUDART, CUDA Driver Version = 5.5, CUDA Runtime Version = 5.0, NumDevs = 1, Device0 = Quadro FX 770M