Bug on multi-dim array element accessing

Nov 5, 2011 at 10:58 AM
Edited Nov 5, 2011 at 2:48 PM

Cudafy is a great library, I like it very much!

I have written a test program, but it seems that cudafy translator has bug on multi-dim array element accessing, and an optional solution is proposed, it will be great to fix the bug in further version.

the C# source code is shown below:

 [Cudafy]

 public static void Test(GThread thread, int[] data)

 {

 int sum = 0;

 for (int dy = -1; dy < 3; dy++)

for (int dx = -1; dx < 3; dx++)

{

sum += kernel[(dy + 1), (dx + 1)];

}

data[0] = sum; 

}

[Cudafy]

public static void Test2(GThread thread, int[, ] data)

{

int sum = 0;

for (int dy = -1; dy < data.GetLength(0) - 1; dy++)

for (int dx = -1; dx < data.GetLength(1) - 1; dx++)

{

sum += data[(dy + 1), (dx + 1)];  <= this line of code is not translated correctly, please refer to following .cu code

}

}

 

The cudafy translator translated the code to the following .cu code:

extern "C" __global__ void Test(int* data, int dataLen0);

extern "C" __global__ void Test2(int* data, int dataLen0, int dataLen1);
__constant__ int kernel[4 * 4];

#define kernelLen0 4

 #define kernelLen1 4

extern "C" __global__ void Test(int* data, int dataLen0)

{

int num = 0;

for (int i = -1; i < 3; i++) {

for (int j = -1; j < 3; j++) {

num += kernel[i + 1 * kernelLen1 +  j + 1];  <== here the code should be num += kernel[(i + 1 * kernelLen1) + ( j + 1)];

}

}

data[0] = num;

}

extern "C" __global__ void Test2(int* data, int dataLen0, int dataLen1)

{

int num = 0;

for (int i = -1; i < dataLen0 - 1; i++){

for (int j = -1; j < dataLen1 - 1; j++) {

num += data[i + 1 * dataLen1 +  j + 1]; <== here the code should be num += kernel[(i + 1 * dataLen1) + ( j + 1)];

}

}

}

One way to fix this problem is to modify CUDAOutputVisitor.cs line 354:

 formatter.WriteToken("(");   // add left bracket

node.AcceptVisitor(this, null);

formatter.WriteToken(")");  // add right bracket

Coordinator
Nov 8, 2011 at 6:24 PM

Thank you very much for reporting this and suggesting the fix.  Much appreciated!

Nick

Coordinator
Nov 9, 2011 at 7:48 AM
Edited Dec 21, 2011 at 6:35 AM
Fix has been committed to repository. However in testing an anomaly was initially
seen where indexing such as [x+1,y+1] generated code with x and y swapped. 
This was subsequently tested in ILSpy 1.0.0.737 which showed same behaviour. 
ILSpy 1.0.0.1000 did not. Subsequent runs of CUDAfy and ILSpy 1.0.0.737 could
no longer reproduce the problem.  
This suggests a minor difference in the IL generated by .NET, 
one that resulted in ILSpy handling it correctly, the other not.
Please be aware of this possible issue and report any findings.
For those of you using not building the sources yourself and 
requiring such indexing with CUDAfy 1.5, or wanting to avoid 
the anomaly described above, here's the workaround:
int x1 = x + 1;
int y1 = y + 1;
num += data[x1, y1];