Cudafy and texture memory

Mar 25, 2012 at 9:50 PM

Hi, 

I have a simple question to which I can't find any sollution. 
I have a kernel write in CUDA C and I figure out how to call it from .net with "DummyFunction" attribute, but How can I copy .net array to texture with Cudafy library?

Could someone paste some code?

Coordinator
Mar 27, 2012 at 7:56 AM

Texture memory is not supported in CUDAfy.  There are of course situations where it is useful but we really want a clean way of doing it.

You can of course access it via the CUDA.NET layer.  Cast your GPGPU object to CudaGPU and then cast the CudaDotNet property.

Mar 27, 2012 at 2:48 PM
Edited Mar 27, 2012 at 2:50 PM

Thanks for your reply.

I did as you say, below sample code:

GPGPU gpu = CudafyHost.GetDevice(CudafyModes.Target);
CudafyModule module = CudafyModule.TryDeserialize("simple_kernel");

if (module == null || !module.TryVerifyChecksums())
{
        module = CudafyTranslator.Cudafy(typeof(Program));
        module.Serialize();
}
gpu.LoadModule(module);

CUDA cuGPU = (CUDA) ((CudaGPU)gpu).CudaDotNet;
var cuTex=cuGPU.GetModuleTexture("simpleTexRef");

var memPtr = cuGPU.CopyHostToDevice(hostData);
cuGPU.SetTextureAddress(cuTex, memPtr, (uint)(sizeof(float) * hostData.Length));

Coordinator
Mar 29, 2012 at 7:57 AM

This is interesting. Can you post some more of your code to show how you use the texture memory? Maybe then we can see how best to access it from CUDAfy in a clean way. Thanks!

Mar 29, 2012 at 10:12 AM

I have own custom CUDA kernel written in CUDA C so  I've only use cudafy to call the kernel with use of "DummyFunctions". 

The simple kernel file "simple_kernel_tex.cu"

texture<float,1,cudaReadModeElementType> simpleTexRef;

extern "C" __global__ void simple_kernel_tex(float * a,int N)
{

	int idx = blockIdx.x*blockDim.x + threadIdx.x;
	
	if (idx<N) a[idx] = a[idx]+tex1Dfetch(simpleTexRef,idx);
}

The full c# program

class Program
    {
        static void Main(string[] args)
        {

            CudafyModes.Target = eGPUType.Cuda;

            GPGPU gpu = CudafyHost.GetDevice(CudafyModes.Target);
            
            int N = 4 * 1024;
            int blockSize=64;
            int gridSize = N/blockSize;
            float[] a = new float[N];
            float[] result = new float[N];
            float[] dev_a = gpu.Allocate<float>(N);

            for (int i = 0; i < N; i++)
            {
                a[i] = i;
            }
            gpu.CopyToDevice(a, dev_a);
            
            module = CudafyModule.TryDeserialize("simple_kernel_tex");
            if (module == null || !module.TryVerifyChecksums())
            {
                module = CudafyTranslator.Cudafy(typeof(Program));
                module.Serialize();
            }
            gpu.LoadModule(module);

            CUDA cuGPU = (CUDA) ((CudaGPU)gpu).CudaDotNet;
            var cuTex=cuGPU.GetModuleTexture("simpleTexRef");

            var memPtr = cuGPU.CopyHostToDevice(a);
            cuGPU.SetTextureAddress(cuTex, memPtr, (uint)(sizeof(float) * a.Length));
         
            gpu.Launch(gridSize, blockSize, "simple_kernel_tex", dev_a, N);
            gpu.CopyFromDevice(dev_a, result);

            for (int i = 0; i < result.Length; i++)
            {
                if (i < 10)
                {
                    if (result[i] != (float)(i + i))
                    {
                        Console.WriteLine("error at {0} position, actual {1} but obtained {2}", i, i + 1, result[i]);
                    }
                    else
                        Console.WriteLine("{0}-> actual {1}, obtained {2}", i, i + i, result[i]);
                }
                else
                    break;
            }
		}

        [CudafyDummy]
        public static void simple_kernel_tex(GThread thread, float[] a, int N)
        {
        }
    }
}

Dec 31, 2014 at 2:04 PM
Edited Dec 31, 2014 at 2:10 PM
Hey, that is really neat what you did with Textures in Cudafy.

https://cudafy.codeplex.com/discussions/576274

I follow the logic, but I have never done mixed C and C#. I am a pure C# programmer & .NET (as well as 45 other languages, assmebler, OS) But it has been 25 years since I did straight C.

So can you clue me in? I am using VS2013. How do I mix C and C# in a project? I see the dummy function and I assume there is some way in the linker to overload the C# dummy with the extern C? But I don't grok this? Do I need a separate C project, create a DLL and then mix them, or is there something simpler? (I have seen mentions of P/Invoke, but I have not dug into that).

Is it some variant on this (Mixing C + + and C#) in VS Studio: http://objectiveseesharp.wordpress.com/2012/07/20/mixing-c-and-c/
Jan 1, 2015 at 3:31 PM
Hi ygutfreund
Nothing could be easier than using dummy functions in cudafy. They are indeed very useful.
Just create a .c file within your project whose name matches that of the function and have it always copy to the output folder. Within, write the function declaration and body in straight c.
In c# create an empty kernel with same signature, and mark it with the dummy attribute.
I think there are several examples in this forum
Jan 1, 2015 at 3:40 PM
Edited Jan 1, 2015 at 4:22 PM
I did read all the examples, and saw the Dummy example (also in the manual) but I did not see the build process. Thanks greatly. BTW i think you mean "*.cu" file not .c file. This will make my team members happy. (and also help in writing 2D Image Convolutions with textures). Here is the discussion article that tells the process, in case others want to find it: https://cudafy.codeplex.com/discussions/479025