Probleme mit Cuda

letavino

Guten Tag!

Ich möchte Cuda in einem bestehenden OpenGL Projekt benutzen,
da sich das Problem aber lediglich auf Cuda bezieht (und ich kein passenderes Unterforum gefunden habe) hoffe ich, ihr verschiebt meinen Thread, sollte er irgendwo besser passen.

Zum Problem:
Ich versuche ein simples Beispiel-Programm zum Addieren von Zahlen mit Cuda zu schreiben, bekomme aber einen Fehler, dessen Ursache ich nicht herausfinden kann.

Der Befehl "cuMemcpyDtoH", also das Kopieren des Speichers von GPU zur CPU schlägt fehl.
Meine Fehlerausgabe sieht folgendermaßen aus:
"Cuda Driver API error = 0700 [...]
String: an illegal memory access was encountered"

Eine Java Implementierung mit Cuda läuft dagegen einwandfrei,
ich kann aber nicht ausschließen, dass das Einbinden in Visual Studio 2012 evtl. nicht korrekt war. Oder kann es andere Ursachen haben, die nichts mit dem Code direkt zutun haben? Ich bin für jegliche Hilfe dankbar!

Sonstige Infos: Windows 10 64bit, gtx 860m, CUDA Version 7.5.18

Hier der Code:

#include <stdio.h>
#include <stdlib.h>

#include <cuda.h>
#include <builtin_types.h>

#define N 100
#define checkCudaErrors(err)  __checkCudaErrors (err, __FILE__, __LINE__)

inline void __checkCudaErrors( CUresult err, const char *file, const int line )
{
    if( CUDA_SUCCESS != err) {
        fprintf(stderr,
                "CUDA Driver API error = %04d from file <%s>, line %i.\n",
                err, file, line );

		const char* str = "";
		const char** pStr = &str;
		cuGetErrorString(err, pStr);
		fprintf(stderr," String: %s",str);
		exit(-1);
	}
}

CUdevice   device;
CUcontext  context;
CUmodule   module;
CUfunction function;
size_t     totalGlobalMem;

char       *module_file = (char*) "matSumKernel.ptx";
char       *kernel_name = (char*) "matSum";

int main(int argc, char **argv)
{
    int a_intArray[N], b_intArray[N], c_intArray[N];
    CUdeviceptr d_a_CUdeviceptr, d_b_CUdeviceptr, d_c_CUdeviceptr;

    // initialize host arrays
    for (int i = 0; i < N; i++) {
        a_intArray[i] = 5;
        b_intArray[i] = 7;
    }
    int deviceCount = 0;

    checkCudaErrors( cuInit(0) );
    checkCudaErrors( cuDeviceGetCount(&deviceCount) );

    if (deviceCount == 0) {
        fprintf(stderr, "Error: no devices supporting CUDA\n");
        exit(-2);
    }
    checkCudaErrors( cuCtxCreate(&context, 0, device) );
    checkCudaErrors( cuModuleLoad(&module, module_file) );
    checkCudaErrors( cuModuleGetFunction(&function, module, kernel_name) );

    // allocate memory
    checkCudaErrors( cuMemAlloc(&d_a_CUdeviceptr, sizeof(int) * N) );
    checkCudaErrors( cuMemAlloc(&d_b_CUdeviceptr, sizeof(int) * N) );
    checkCudaErrors( cuMemAlloc(&d_c_CUdeviceptr, sizeof(int) * N) );

    // copy arrays to device
    checkCudaErrors( cuMemcpyHtoD(d_a_CUdeviceptr, a_intArray, sizeof(int) * N) );
    checkCudaErrors( cuMemcpyHtoD(d_b_CUdeviceptr, b_intArray, sizeof(int) * N) );

    // run
    printf("# Running the kernel...\n");

    void *args[3] = { &d_a_CUdeviceptr, &d_b_CUdeviceptr, &d_c_CUdeviceptr };

    checkCudaErrors( cuLaunchKernel(function, 
	N, 1, 1,			// Nx1x1 blocks
        1, 1, 1,            // 1x1x1 threads
        0, 0, args, 0) );
    printf("# Kernel complete.\n");

    // copy results to host and report
    //-------------NÄCHSTE ZEILE BRINGT FEHLER-------------------------
    checkCudaErrors( cuMemcpyDtoH(c_intArray, d_c_CUdeviceptr, sizeof(int) * N) );
   //[...] wird nicht mehr erreicht	
}

Cuda Kernel:

extern "C" __global__ void matSum(int *a, int *b, int *c)
{
    int index = threadIdx.x + blockIdx.x * blockDim.x;
    if (index < 100)
        c[index] = a[index] + b[index];
}

megaweber

Hi, in welcher Zeile taucht der Fehler auf? In deiner Beschreibung sprichst du von Device2Host im Code von Host2Device, sicher dass Du den korrekten Code gepostet hast? Mit der Cuda Driver API kenne ich nicht so gut aus, hast du einen Grund diese statt der üblicheren Runtime API zu verwenden?

letavino

megaweber schrieb:

Hi, in welcher Zeile taucht der Fehler auf? In deiner Beschreibung sprichst du von Device2Host im Code von Host2Device, sicher dass Du den korrekten Code gepostet hast? Mit der Cuda Driver API kenne ich nicht so gut aus, hast du einen Grund diese statt der üblicheren Runtime API zu verwenden?

Hallo,
der Fehler taucht ganz unten in Zeile 79 auf, beim zurückkopieren vom Device zum Host.
Die Driver Api habe ich vor allem genommen, da ich eine bestehende Java Implementierung mit JCuda in c/c++ umschreiben möchte. JCuda unterstützt nur die Driver Api und ich dachte, so ist der Umstieg leichter.
Außerdem kam meine Umgebung beim testen mit den <<< und >>> nicht klar, da habe ich mich dann auch nicht weiter mit aufgehalten.

dot

Ich würde mal vermuten, dass das Problem nicht beim cuMemcpyDtoH() an sich liegt, sondern dass dein Kernel bei der Ausführung einen out-of-bounds Access macht. Der wird auf CPU Seite lediglich erst beim nächsten Sync mit der GPU bemerkt, was in deinem Fall implizit im cuMemcpyDtoH() nach dem Launch passiert. Wenn du vor dem cuMemcpyDtoH() erst ein explizites cuCtxSynchronize() machst, wirst du feststellen, dass der Fehler auf einmal dort auftritt. Es ist generell so, dass CUDA API Calls mit ihrem Rückgabewert über den Fehlschlag früherer Operationen berichten können...

Lass dein Ding mal durch Nsight mit aktiviertem Memchecker laufen, du hast ziemlich sicher einen Bug in deinem Kernel...

megaweber schrieb:

Mit der Cuda Driver API kenne ich nicht so gut aus, hast du einen Grund diese statt der üblicheren Runtime API zu verwenden?

letavino schrieb:

Die Driver Api habe ich vor allem genommen, da ich eine bestehende Java Implementierung mit JCuda in c/c++ umschreiben möchte. JCuda unterstützt nur die Driver Api und ich dachte, so ist der Umstieg leichter.
Außerdem kam meine Umgebung beim testen mit den <<< und >>> nicht klar, da habe ich mich dann auch nicht weiter mit aufgehalten.

Glückliche Entscheidung; Ich will gar nicht wissen, wer auf die glorreiche Idee kam, dass man Code, der durch zwei grundverschiedene Compiler läuft und dabei jeweils für zwei grundverschiedene Architekturen übersetzt wird, am besten in ein gemeinsames File packt, nur damit man es durch die schwarze Magie eines dritten Frontend, das die Dinger am Ende in mehrere Files splitted, mit aller Kraft so aussehen lassen kann, als ob alles irgendwie das selbe wäre...Hab längst aufgehört zu zählen, wie oft wieder mal jemand bei mir in der Tür stand, weil diese Wundermaschine wieder mal unbemerkt eine seiner Extremitäten gefressen hatte. Der CUDA Runtime API Kram ist der gröbste Unfug, ich kann nur jedem ans Herz legen, sich soweit davon entfernt zu halten, wie nur irgendwie möglich...

Selbst in deinem Driver API Code machst du etwas, das streng genommen problematisch sein kann: sizeof(int) muss aufseiten der CPU nicht notwendigerweise das selbe sein, wie aufseiten der GPU. Ich persönlich würde im Hostcode besser explizit mit std::int32 arbeiten. Genaugenommen müsste man sich ja selbst auch noch über so Dinge wie Endianess und Integerrepresentation Gedanken machen...

Es sei auch darauf hingewiesen, dass im globalen Namespace alle Namen, die zwei Underscores enthalten, reservierte Namen sind. Sowas wie dein __checkCudaErrors() bringt dir streng genommen Undefined Behavior...

Und N Blöcke von 1x1x1 Threads sind so ziemlich das ineffizienteste, was man auf der GPU nur machen kann...

rapso

laeuft das ganze wenn du zeile 71 bis 74 auskommentierst?

letavino

@rapso:
Ja, wenn ich die Funktion auskommentiere, läuft das Programm weiter.
Auch ein Hinweis darauf, dass der Kernel fehlerhaft ist?

@dot:
Vielen Dank erstmal für deinen ausführlichen und gut erklärenden Beitrag!
Hättest du eine Vermutung, wie es bei den wenigen Zeilen des Kernels zu einem Fehler kommen kann?
(Der überschauliche Kernelcode ist unterhalb des restlichen Quellcodes.)
Nsight werde ich mir mal besorgen.
Vielen Dank auch für die weiteren Hinweise zu meinem Code!
Es handelt sich hier aber erstmal um einen kleinen Test, um Cuda überhaupt zum Laufen zu bringen, da ist es nicht so schlimm, wenn die hundertmalige Berechnung von 5+7 nicht ganz so effizient ist. (Wird natürlich im richtigen Programm dann angepasst).

letavino

Fehler gefunden!

Dank euch habe ich an der richtige Stelle gesucht, dem Kernel.
Und dank deiner Anmerkung über die mögliche unterschiedliche Definition von "int" auf der CPU und GPU, fand ich den Fehler schnell bei einem Blick in die ptx Datei.

[...]
.visible .entry matSum(
	.param .u64 matSum_param_0,
[...]

Die Werte der Arrays werden von der GPU als 64bit Variablen interpretiert, von der CPU aber als 32bit.
Eine kleine Anpassung beim NVCC (-m 32) und schon läuft das Programm einwandfrei. (Müsste genauso beides auf 64bit laufen.)

Vielen Dank! Die Hilfe und die Fragen haben mich genau zu dem bestehenden Problem geführt!

rapso

ueber sizeof(int) muss man sich keine sorgen machen, du wirst keinen compiler finden der cuda unterstuetzt (NV-sdk) und dabei was anderes als cuda ausspuckt.

aber man muss natuerlich dieselbe platform compilieren (cpu+gpu). dabei ist der cpu-seitige code schneller mit 64bit, cuda ist schneller mit 32bit.

dot

Nur aus Interesse: Was für einen Compiler/Toolchain verwendest du bzw. wie kompilierst du den CUDA Code? Das Problem hier war nicht das sizeof(int) (wie rapso schon gesagt hat, ist das momentan nur ein theoretisches Problem, da es in der Praxis kein von CUDA supportetes System gibt, wo eine entsprechende Konfiguration anzutreffen wäre, imo sollte man sich dessen aber zumindest bewusst sein), sondern dass deine CPU Seite wohl für 32-Bit kompiliert wurde und CUdeviceptr daher auch nur 32-Bit war, dein GPU Code aber 64-Bit Pointer erwartet hat. Effektiv wurden dadurch wohl auf GPU Seite zwei Pointerargumente zu einem zusammengeklebt und der resultierende Pointer zeigte natürlich ins Nirvana...