diff --git a/GPU.md b/GPU.md
index 7b57db4..d748c4a 100644
--- a/GPU.md
+++ b/GPU.md
@@ -488,7 +488,7 @@ int main() {
 }
 ```
 
-## 1.8.使用`atomicAdd`避免数据竞争.
+## 1.8.使用`atomicAdd`避免数据竞争
 
 ```c++
 #include <iostream>
@@ -524,6 +524,31 @@ int main() {
 }
 ```
 
+## 1.9.使用`__device__`函数
+
+```c++
+#include <cstdio>
+
+// 声明的__device__函数只能在GPU上被调用.
+__device__ const char* GetString() {
+    return "Hello, World\n";
+}
+
+__global__ void PrintString() {
+    // 核函数只能调用__device__函数.
+    const char *string = GetString();
+
+    printf("%s", string);
+}
+
+int main() {
+    PrintString<<<1, 1>>>();
+    cudaDeviceSynchronize();
+
+    return 0;
+}
+```
+
 # 2.cuBLAS
 
 ## 2.1.cublas\<t\>gemm()