CUDA.NET & Textures

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi,

I am trying to get CUDA.NET work together with textures. But no success for now. Here is the small test code. But it returns only zeros in the fout :(

C# code:

[codebox]using System;

using System.Collections.Generic;

using System.Linq;

using System.Text;

using GASS.CUDA;

using GASS.CUDA.Types;

namespace CUDA_Texture

{

class Program

{

	static void Main(string[] args)

	{

		CUDA cuda = new CUDA(0, true);

		CUmodule module = cuda.LoadModule("texture.cubin");

		CUfunction cf = cuda.GetModuleFunction(module, "testTexture");

		CUtexref textRef = cuda.GetModuleTexture(module, "ttex");

		int n = 32;

		uint[] a = new uint[n];

		for (int i = 0; i < n; i++)

		{

			a[i] = (uint)i;

		}

		CUdeviceptr d_a = cuda.CopyHostToDevice<uint>(a);

		cuda.SetTextureAddressMode(textRef, 0, CUAddressMode.Clamp);

		cuda.SetTextureFilterMode(textRef, CUFilterMode.Point);

		cuda.SetTextureFormat(textRef, CUArrayFormat.UnsignedInt32, 1);

		cuda.SetTextureFlags(textRef, CUDADriver.CU_TRSF_READ_AS_INTEGER);

		cuda.SetTextureAddress(textRef, d_a, (uint)(n) * sizeof(uint));

		uint[] fout = new uint[n];

		CUdeviceptr d_out = cuda.Allocate<uint>(fout);

		cuda.SetFunctionBlockShape(cf, n, 1, 1);

		cuda.SetParameter(cf, 0, d_out.Pointer);

		cuda.SetParameterSize(cf, sizeof(uint));

		cuda.Launch(cf, 1, 1);

		cuda.CopyDeviceToHost<uint>(d_out, fout);

	}

}

}[/codebox]

C kernel (compiled into ‘texture.cubin’)

[codebox]typedef unsigned int uint;

texture<unsigned int, 1, cudaReadModeElementType> ttex;

extern “C”

global void testTexture(uint* out)

{

uint index = threadIdx.x;

int i = (int)index;

out[index] = tex1Dfetch(ttex, i);

}[/codebox]

Anybody could help me?

Thanks

EDIT: I am usung CUDA 2.2, CUDA.NET 2.0, GeForce 9800GTX, WinXP, VS2008

2

Hi Dusan,

When you want a texture to be accessible to a specific kernel or function, you need to mark it as a parameter too.

The same way that you specify parameters in order for the kernel, use the SetParameter function that accepts a texture reference.

This will enable the texture data be available inside the function.

Besides that, the code you specified is OK and the steps you perform are correct.

Best,

Moti.

3

It was so easy :)
It works now.
Thanks.

4

Is it possible to post some working code here I’m trying to do almost exactly the same thing.

However I’m getting issues on the kernel side using textures as parameters.

When I try to compile something as simple as

include <vector_types.h>

extern “C” global void partialFFT_withNorm(texture<int,1,cudaReadModeElementType> &colD_)

{						

int test = tex1Dfetch(colD_,4);

}

It fails to output the .cubin file and gives this error.

It’s a weird one, I could have swore it was working before lunch?

1>Signal: caught in Code_Expansion phase.

1>(0): Error: Signal caught in phase Code_Expansion – processing aborted

1>nvopencc ERROR: C:\CUDA\bin/…/open64/lib//be.exe returned non-zero status 3

5

Hi enigma015,

I was on holiday last week…

My code that works is bellow. I’ve just added one row as moti.bot said (marked red). You should add the parameter to cuda function (use SetParameter), not add an argument to kernel. Here is the code:

[codebox]using System;

using System.Collections.Generic;

using System.Linq;

using System.Text;

using GASS.CUDA;

using GASS.CUDA.Types;

namespace CUDA_Texture

{

class Program

{

	static void Main(string[] args)

	{

		CUDA cuda = new CUDA(0, true);

		CUmodule module = cuda.LoadModule("texture.cubin");

		CUfunction cf = cuda.GetModuleFunction(module, "testTexture");

		CUtexref textRef = cuda.GetModuleTexture(module, "ttex");

		int n = 32;

		uint[] a = new uint[n];

		for (int i = 0; i < n; i++)

		{

			a[i] = (uint)i;

		}

		CUdeviceptr d_a = cuda.CopyHostToDevice<uint>(a);

		cuda.SetTextureAddressMode(textRef, 0, CUAddressMode.Clamp);

		cuda.SetTextureFilterMode(textRef, CUFilterMode.Point);

		cuda.SetTextureFormat(textRef, CUArrayFormat.UnsignedInt32, 1);

		cuda.SetTextureFlags(textRef, CUDADriver.CU_TRSF_READ_AS_INTEGER);

		cuda.SetTextureAddress(textRef, d_a, (uint)(n) * sizeof(uint));

		uint[] fout = new uint[n];

		CUdeviceptr d_out = cuda.Allocate<uint>(fout);

		cuda.SetFunctionBlockShape(cf, n, 1, 1);

		cuda.SetParameter(cf, 0, d_out.Pointer);

		cuda.SetParameterSize(cf, sizeof(uint));

		<b>cuda.SetParameter(cf, textRef);</b>

		cuda.Launch(cf, 1, 1);

		cuda.CopyDeviceToHost<uint>(d_out, fout);

	}

}

}

[/codebox]

The *.cu file is the same as it was before:

[codebox]typedef unsigned int uint;

texture<unsigned int, 1, cudaReadModeElementType> ttex;

extern “C”

global void testTexture(uint* out)

{

uint index = threadIdx.x;

int i = (int)index;

out[index] = tex1Dfetch(ttex, i);

}[/codebox]

Hope it will help you.

Now I am trying to make 3D texture works…