HI,
I made an OpenACC code where most of variables are double type.
Before, I used NVIDIA Geforce GPU. But, in order to reduce the calculation time, I bought AMD FirePro W9100.
The code compiled without problem with compiler option -ta=radeon:hawaii.
(I’m using PGI Community Edition Version 16.10. After installing AMD driver, I reinstalled PGI Community Edition.)
However, when I run it, the data clause does not work even though pgaccelinfo says that I have AMD FIrePro W9100.
(In the below sample code, in the output file, x and y are arbitrary numbers. Even if I use -ta=radeon option, the problem is not solved.)
Compiling the same code with compiler option -ta=multicore or -ta= tesla:cuda8.0 (when I used Geforce GPU which is compatible with cuda8.0) does not bring any problem and data clause works perfectly fine.
Could you give me some help so that the below code works with AMD GPU?
If OpenACC is not compatible with AMD GPU, should I have to use OpenCL?
(If so, is there any simple way to convert OpenACC code to OpenCL code? I’m afraid that it would take too much time to run OpenCL)
Thanks in advance.
Here are my test code. It consists of main code file, header file. If one compiles with “pgcc -acc -ta=radeon:hawaii -Minfo=accel amdtest.c -o amdtest” and the input file is “testinput”, then one should type “./amdtest -p testinput” in order to execute the program.
#include "test.h"
int main(int argc, char **argv) {
FILE *out;
int optnum, optcheck;
long cnti, cntj;
double *restrict x;
double *restrict y;
if ((argc < 3) || (strcmp(*(argv + 1), "-p") != 0)) {
fprintf(stderr, "Usage: %s -p <parameterfile> \n", *argv);
exit(EXIT_FAILURE);
}
optnum = argc - 1;
printf("Parameter numbers = %d\n", optnum - 1);
for(optcheck = 2; optcheck <= optnum; optcheck++)
{
if (!cfg_init(argv[optcheck])) {
fprintf(stderr, "Wrong input parameter file.\n");
exit(EXIT_FAILURE);
}
readpar(argv[optcheck]);
x = alloc_double_vector(Nx);
y = alloc_double_vector(Ny);
#pragma acc data copyout(x[:Nx], y[:Ny]) pcopyin(Nx, Ny, dx, dy, Nx2, Ny2)
{
#pragma acc kernels loop
for (cnti = 0; cnti < Nx; cnti++) {
x[cnti] = (cnti - Nx2) * dx;
}
#pragma acc kernels loop
for (cntj = 0; cntj < Ny; cntj++) {
y[cntj] = (cntj - Ny2) * dy;
}
}
if (outname != NULL) {
out = fopen(outname, "w");
for (cnti = 0; cnti < Nx; cnti += outstpx) {
for (cntj = 0; cntj < Ny; cntj += outstpy) {
fprintf(out, "%e %e\n", x[cnti], y[cntj]);
}
}
fclose(out);
}
free_double_vector(x);
free_double_vector(y);
}
return(EXIT_SUCCESS);
}
int cfg_init(char *cfg_file) {
FILE *file;
char buf[256];
file = fopen(cfg_file, "r");
if (!file) return 0;
cfg_size = 0;
while (fgets(buf, 256, file) != NULL) {
if (sscanf(buf, "%s = %s", cfg_key[cfg_size], cfg_val[cfg_size]) == 2) cfg_size++;
}
fclose(file);
return cfg_size;
}
char *cfg_read(char *key) {
int i;
for (i = 0; i < cfg_size; i++)
if (!strcmp(key, cfg_key[i])) return cfg_val[i];
return NULL;
}
double *alloc_double_vector(long Nx1) {
double *vector;
if ((vector = (double *)malloc((size_t)(Nx1 * sizeof(double)))) == NULL) {
fprintf(stderr, "Failed to allocate memory for the vector.\n");
exit(EXIT_FAILURE);
}
return vector;
}
void free_double_vector(double *vector) {
free((char *)vector);
}
void readpar(char *cfg_tmp) {
if ((cfg_tmp = cfg_read("NX")) == NULL) {
fprintf(stderr, "NX is not defined in the configuration file.\n");
exit(EXIT_FAILURE);
}
Nx = atol(cfg_tmp);
if(Nx % 2 == 0) {
Nx = Nx + 1;
printf("Nx is changed.\n");
}
Nx2 = Nx / 2;
if ((cfg_tmp = cfg_read("NY")) == NULL) {
fprintf(stderr, "NY is not defined in the configuration file.\n");
exit(EXIT_FAILURE);
}
Ny = atol(cfg_tmp);
if(Ny % 2 == 0) {
Ny = Ny + 1;
printf("Ny is changed.\n");
}
Ny2 = Ny / 2;
if ((cfg_tmp = cfg_read("DX")) == NULL) {
fprintf(stderr, "DX is not defined in the configuration file.\n");
exit(EXIT_FAILURE);
}
dx = atof(cfg_tmp);
if ((cfg_tmp = cfg_read("DY")) == NULL) {
fprintf(stderr, "DY is not defined in the configuration file.\n");
exit(EXIT_FAILURE);
}
dy = atof(cfg_tmp);
outname = cfg_read("OUTNAME");
if ((outname != NULL)) {
if ((cfg_tmp = cfg_read("OUTSTPX")) == NULL) {
fprintf(stderr, "OUTSTPX is not defined in the configuration file.\n");
exit(EXIT_FAILURE);
}
outstpx = atol(cfg_tmp);
if ((cfg_tmp = cfg_read("OUTSTPY")) == NULL) {
fprintf(stderr, "OUTSTPY is not defined in the configuration file.\n");
exit(EXIT_FAILURE);
}
outstpy = atol(cfg_tmp);
}
return;
}
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <math.h>
long Nx, Ny;
long Nx2, Ny2;
double dx, dy;
char *outname;
long outstpx, outstpy;
void readpar(char *);
double *alloc_double_vector(long);
void free_double_vector(double *);
int cfg_init(char *);
char *cfg_read(char *);
int cfg_size;
char cfg_key[256][256], cfg_val[256][256];
# Number of discretization points in the x axis.
# Type: long
NX = 101
# Number of discretization points in the y axis.
# Type: long
NY = 101
# Spatial discretization step in the x axis.
# Type: double
DX = 0.01
# Spatial discretization step in the y axis.
# Type: double
DY = 0.01
# Output file
# Type: string
OUTNAME = testoutput.txt
# Discretization step in the x-direction used to save x.
# Type: long
OUTSTPX = 1
# Discretization step in the y-direction used to save y.
# Type: long
OUTSTPY = 1