Skip to content

Commit fe53cfa

Browse files
harrismharrism
committed
Adding examples for Unified Memory CUDACast
1 parent 692f464 commit fe53cfa

6 files changed

Lines changed: 306 additions & 0 deletions

File tree

‎posts/unified-memory/.gitignore‎

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,4 @@
1+
dataElem
2+
dataElem_um
3+
dataElem_um_c++_1
4+
dataElem_um_c++_2

‎posts/unified-memory/Makefile‎

Lines changed: 13 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,13 @@
1+
CUDA_ARCH_FLAGS := -arch=sm_35
2+
CC_FLAGS += $(CUDA_ARCH_FLAGS)
3+
4+
EXE = dataElem dataElem_um dataElem_um_c++_1 dataElem_um_c++_2
5+
6+
all: $(EXE)
7+
8+
% : %.cu
9+
nvcc $< $(CC_FLAGS) $(LIB_FLAGS) -o $@
10+
11+
clean:
12+
rm -f $(EXE)
13+

‎posts/unified-memory/dataElem.cu‎

Lines changed: 60 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,60 @@
1+
#include <string.h>
2+
#include <stdio.h>
3+
4+
struct DataElement
5+
{
6+
char *name;
7+
int value;
8+
};
9+
10+
__global__
11+
void Kernel(DataElement *elem) {
12+
printf("On device: name=%s, value=%d\n", elem->name, elem->value);
13+
14+
elem->name[0] = 'd';
15+
elem->value++;
16+
}
17+
18+
void launch(DataElement *elem) {
19+
DataElement *d_elem;
20+
char *d_name;
21+
22+
int namelen = strlen(elem->name) + 1;
23+
24+
// Allocate storage for struct and text
25+
cudaMalloc(&d_elem, sizeof(DataElement));
26+
cudaMalloc(&d_name, namelen);
27+
28+
// Copy up each piece separately, including new “text” pointer value
29+
cudaMemcpy(d_elem, elem, sizeof(DataElement), cudaMemcpyHostToDevice);
30+
cudaMemcpy(d_name, elem->name, namelen, cudaMemcpyHostToDevice);
31+
cudaMemcpy(&(d_elem->name), &d_name, sizeof(char*), cudaMemcpyHostToDevice);
32+
33+
// Finally we can launch our kernel, but CPU & GPU use different copies of “elem”
34+
Kernel<<< 1, 1 >>>(d_elem);
35+
36+
cudaMemcpy(&(elem->value), &(d_elem->value), sizeof(int), cudaMemcpyDeviceToHost);
37+
cudaMemcpy(elem->name, d_name, namelen, cudaMemcpyDeviceToHost);
38+
39+
cudaFree(d_name);
40+
cudaFree(d_elem);
41+
}
42+
43+
int main(void)
44+
{
45+
DataElement *e;
46+
e = (DataElement*)malloc(sizeof(DataElement));
47+
48+
e->value = 10;
49+
e->name = (char*)malloc(sizeof(char) * strlen("hello") );
50+
strcpy(e->name, "hello");
51+
52+
launch(e);
53+
54+
printf("On host: name=%s, value=%d\n", e->name, e->value);
55+
56+
free(e->name);
57+
free(e);
58+
59+
cudaDeviceReset();
60+
}

‎posts/unified-memory/dataElem_um.cu‎

Lines changed: 40 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,40 @@
1+
#include <string.h>
2+
#include <stdio.h>
3+
4+
struct DataElement
5+
{
6+
char *name;
7+
int value;
8+
};
9+
10+
__global__
11+
void Kernel(DataElement *elem) {
12+
printf("On device: name=%s, value=%d\n", elem->name, elem->value);
13+
14+
elem->name[0] = 'd';
15+
elem->value++;
16+
}
17+
18+
void launch(DataElement *elem) {
19+
Kernel<<< 1, 1 >>>(elem);
20+
cudaDeviceSynchronize();
21+
}
22+
23+
int main(void)
24+
{
25+
DataElement *e;
26+
cudaMallocManaged((void**)&e, sizeof(DataElement));
27+
28+
e->value = 10;
29+
cudaMallocManaged((void**)&(e->name), sizeof(char) * strlen("hello") );
30+
strcpy(e->name, "hello");
31+
32+
launch(e);
33+
34+
printf("On host: name=%s, value=%d\n", e->name, e->value);
35+
36+
cudaFree(e->name);
37+
cudaFree(e);
38+
39+
cudaDeviceReset();
40+
}

‎posts/unified-memory/dataElem_um_c++_1.cu‎

Lines changed: 54 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,54 @@
1+
#include <string.h>
2+
#include <stdio.h>
3+
4+
class Managed
5+
{
6+
public:
7+
void *operator new(size_t len) {
8+
void *ptr;
9+
cudaMallocManaged(&ptr, len);
10+
return ptr;
11+
}
12+
13+
void operator delete(void *ptr) {
14+
cudaFree(ptr);
15+
}
16+
};
17+
18+
struct DataElement : public Managed
19+
{
20+
char *name;
21+
int value;
22+
};
23+
24+
__global__
25+
void Kernel(DataElement *elem) {
26+
printf("On device: name=%s, value=%d\n", elem->name, elem->value);
27+
28+
elem->name[0] = 'd';
29+
elem->value++;
30+
}
31+
32+
void launch(DataElement *elem) {
33+
Kernel<<< 1, 1 >>>(elem);
34+
cudaDeviceSynchronize();
35+
}
36+
37+
int main(void)
38+
{
39+
DataElement *e = new DataElement;
40+
41+
e->value = 10;
42+
cudaMallocManaged((void**)&(e->name), sizeof(char) * strlen("hello") );
43+
strcpy(e->name, "hello");
44+
45+
launch(e);
46+
47+
printf("On host: name=%s, value=%d\n", e->name, e->value);
48+
49+
cudaFree(e->name);
50+
delete e;
51+
52+
cudaDeviceReset();
53+
}
54+

‎posts/unified-memory/dataElem_um_c++_2.cu‎

Lines changed: 135 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,135 @@
1+
#include <string.h>
2+
#include <stdio.h>
3+
4+
// Managed Base Class -- inherit from this to automatically
5+
// allocate objects in Unified Memory
6+
class Managed
7+
{
8+
public:
9+
void *operator new(size_t len) {
10+
void *ptr;
11+
cudaMallocManaged(&ptr, len);
12+
return ptr;
13+
}
14+
15+
void operator delete(void *ptr) {
16+
cudaFree(ptr);
17+
}
18+
};
19+
20+
// String Class for Managed Memory
21+
class String : public Managed
22+
{
23+
int length;
24+
char *data;
25+
public:
26+
String() : length(0), data(0) {}
27+
28+
// Constructor for C-string initializer
29+
String(const char *s) {
30+
_realloc(strlen(s));
31+
memcpy(data, s, length);
32+
}
33+
34+
// Copy constructor
35+
String(const String& s) {
36+
_realloc(s.length);
37+
memcpy(data, s.data, length);
38+
}
39+
40+
~String() { cudaFree(data); }
41+
42+
// Assignment operator
43+
String& operator=(const char* s) {
44+
_realloc(strlen(s));
45+
memcpy(data, s, length);
46+
return *this;
47+
}
48+
49+
// Element access (from host or device)
50+
__host__ __device__
51+
char& operator[](int pos) { return data[pos]; }
52+
53+
// C-string access
54+
__host__ __device__
55+
const char* c_str() const { return data; }
56+
57+
private:
58+
void _realloc(int len) {
59+
cudaFree(data);
60+
length = len;
61+
cudaMallocManaged(&data, length);
62+
}
63+
};
64+
65+
66+
struct DataElement : public Managed
67+
{
68+
String name;
69+
int value;
70+
};
71+
72+
__global__
73+
void Kernel_by_pointer(DataElement *elem) {
74+
printf("On device: name=%s, value=%d\n", elem->name.c_str(), elem->value);
75+
76+
elem->name[0] = 'p';
77+
elem->value++;
78+
}
79+
80+
__global__
81+
void Kernel_by_ref(DataElement &elem) {
82+
printf("On device: name=%s, value=%d\n", elem.name.c_str(), elem.value);
83+
84+
elem.name[0] = 'r';
85+
elem.value++;
86+
}
87+
88+
__global__
89+
void Kernel_by_value(DataElement elem) {
90+
printf("On device: name=%s, value=%d\n", elem.name.c_str(), elem.value);
91+
92+
elem.name[0] = 'v';
93+
elem.value++;
94+
}
95+
96+
void launch_by_pointer(DataElement *elem) {
97+
Kernel_by_pointer<<< 1, 1 >>>(elem);
98+
cudaDeviceSynchronize();
99+
}
100+
101+
void launch_by_ref(DataElement &elem) {
102+
Kernel_by_ref<<< 1, 1 >>>(elem);
103+
cudaDeviceSynchronize();
104+
}
105+
106+
void launch_by_value(DataElement elem) {
107+
Kernel_by_value<<< 1, 1 >>>(elem);
108+
cudaDeviceSynchronize();
109+
}
110+
111+
int main(void)
112+
{
113+
DataElement *e = new DataElement;
114+
115+
e->value = 10;
116+
e->name = "hello";
117+
118+
launch_by_pointer(e);
119+
120+
printf("On host (after by-pointer): name=%s, value=%d\n", e->name.c_str(), e->value);
121+
122+
launch_by_ref(*e);
123+
124+
printf("On host (after by-ref): name=%s, value=%d\n", e->name.c_str(), e->value);
125+
126+
launch_by_value(*e);
127+
128+
printf("On host (after by-value): name=%s, value=%d\n", e->name.c_str(), e->value);
129+
130+
delete e;
131+
132+
cudaDeviceReset();
133+
}
134+
135+

0 commit comments

Comments
 (0)