std::map in device code

Accelerated Computing CUDA CUDA Programming and Performance
1

Is anyone aware of an implementation of STL containers that can be used in device code? I was able to find this [url]STL::map - CUDA Programming and Performance - NVIDIA Developer Forums thread on the topic but no actual code.

I’m aware of the implications of doing this on a GPU and I still want to try it. I’m just posting this in case someone has already done it, but I’m assuming that no one has, and will just try to port the implementation in LLVM’s libcxx. The code will go here if I succeed: [url]Google Code Archive - Long-term storage for Google Code Project Hosting.

2

If anyone is interested, I was able to get this to work. It was actually surprisingly easy.

I haven’t tried any series template programming in CUDA since back around 2.3,
and NVCC has gotten so much better since then.

Here’s an example.

#include <archaeopteryx/util/interface/map.h>

#include <cstdio>

__global__ void testMap()
{

	archaeopteryx::util::map<int, int> deviceMap;
	
	deviceMap.insert(archaeopteryx::util::make_pair(0, 10));
	deviceMap.insert(archaeopteryx::util::make_pair(30, 5));
	deviceMap.insert(archaeopteryx::util::make_pair(15, 2));
	
	printf("Map values:\n");

	for(archaeopteryx::util::map<int, int>::iterator i = deviceMap.begin();
		i != deviceMap.end(); ++i)
	{
		printf(" (%d, %d)\n", i->first, i->second);
	}
}

int main(int argc, char** argv)
{
	testMap<<<1,1>>>();
	cudaError_t success = cudaDeviceSynchronize();
	
	if(success != cudaSuccess)
	{
		std::printf("Test Failed\n");
		std::printf(" Message: %s\n", cudaGetErrorString(success));
	}
	else
	{
		std::printf("Test Passed\n");
	}
	
	return 0;
}

normal@diamos:~/checkout/vanaheimr/archaeopteryx$ nvcc example.cu -arch sm_20 -I . -o example
ptxas warning : Stack size for entry function '_Z7testMapv' cannot be statically determined
normal@diamos:~/checkout/vanaheimr/archaeopteryx$ ./example 
Map values:
 (0, 10)
 (15, 2)
 (30, 5)
Test Passed

The code is derived from libcxx and BSD licensed in case someone wants to take it further. There is a
partial implementation of string, vector, algorithms, memory, and all of the other necessary baggage.

edit: the comment box ate a bunch of my code

3

/home/archaeopteryx$ nvcc example.cu -arch sm_61 -I . -I …/ -o example
ptxas warning : Stack size for entry function ‘_Z7testMapv’ cannot be statically determined
./archaeopteryx/util/interface/map.h(510): error: name followed by “::” must be a class or namespace name

1 error detected in the compilation of “/tmp/tmpxft_000498e6_00000000-4_example.cpp4.ii”.

4

I was able to work around this by changing the offending line (510) in map.h from:

const_iterator __next = _VSTD::next(__hint);

to:

const_iterator __next = std::next(__hint);

and also adding:

#include <iterator>

at the top of the file, and adding

-std=c++11

to the compile command line.

_VSTD seems to be something that existed once but may have been removed from language implementations:
https://reviews.llvm.org/D55517
but that is just a guess.