CUDA

什么是CUDA

CUDA是nvida基于自家GPU提供的一套并行計(jì)算框架，通過CUDA可以很方便的進(jìn)行編碼利用GPU強(qiáng)大的并行處理能力完成對(duì)應(yīng)任務(wù)計(jì)算。

基本概念

HOST（主機(jī)部分）：CPU以及CPU所使用的內(nèi)存以及在CPU上執(zhí)行的代碼。

DEVICE（設(shè)備）：GPU以及GPU所使用的顯存以及對(duì)應(yīng)運(yùn)行在GPU上的代碼。

基本邏輯結(jié)構(gòu)如下圖

Kernel(核函數(shù))：GPU上執(zhí)行的函數(shù)被稱為核函數(shù)。

Thread（線程）：GPU所執(zhí)行的最小單元（一般由GPU上的一個(gè)處理核執(zhí)行）受神。

Block（線程塊）：多個(gè)Thread組成的一組線程，Block內(nèi)的線程可以進(jìn)行數(shù)據(jù)共享格侯，Block之間無法進(jìn)行數(shù)據(jù)共享鼻听，不同的block是并行執(zhí)行的，并且不同的block執(zhí)行上沒有先后順序联四。在一個(gè)Block中撑碴，一般最多創(chuàng)建不超過512個(gè)Thread。

Grid（線程格）：有多個(gè)Block組成朝墩，一般一個(gè)Grid內(nèi)的線程數(shù)量不超過65535個(gè)（該Grid內(nèi)所有Block線程之和）醉拓，一般一個(gè)GPU設(shè)備可以生成2到3個(gè)Grid。

Wrap：邏輯上收苏，所有的thread都是并行執(zhí)行亿卤，但從硬件的角度看，并不是所有的thread都能夠在同一時(shí)間執(zhí)行鹿霸，Wrap是GPU調(diào)度和執(zhí)行thread的最基本單元（可理解為GPU硬件執(zhí)行單元）排吴，一般情況下，一個(gè)wrap包含32個(gè)thread懦鼠，這32個(gè)thread執(zhí)行同一條處理指令（數(shù)據(jù)不同）钻哩。此外一個(gè)wrap中被執(zhí)行的線程必然屬于同一個(gè)Block，如果Block中的線程數(shù)目不是wrap大小的整數(shù)倍葛闷，在block被執(zhí)行時(shí)憋槐，也會(huì)被湊夠?yàn)閣rap的整數(shù)倍（多余的線程只是狀態(tài)被設(shè)置為非活動(dòng)狀態(tài)）双藕。

Warp Divergence (Wrap分歧)：CPU擁有非常好的分支預(yù)測(cè)能力淑趾，如果預(yù)測(cè)正確，CPU只會(huì)由很小的消耗忧陪，和CPU相比扣泊，GPU的分支預(yù)測(cè)就非常差近范，同一個(gè)wrap中的thread同時(shí)執(zhí)行相同的指令，如果thread遇到分支控制語句后延蟹，不同thread進(jìn)入不同的分支评矩，就會(huì)導(dǎo)致其余分支都被阻塞，十分影響性能阱飘，這類問題就被稱為Warp Divergence斥杜，Warp Divergence只會(huì)發(fā)生在同一個(gè)wrap中。

???????? SM（Streaming Multiprocessor）：GPU硬件上的一個(gè)概念沥匈，一個(gè)GPU有非常多的SM組成蔗喂，一個(gè)SM內(nèi)部，會(huì)有該SM對(duì)應(yīng)的寄存器高帖，指令cache缰儿，指令buffer，wrap調(diào)度器散址，計(jì)算core等部分乖阵。一般情況下一個(gè)Wrap會(huì)被分配到一個(gè)SM上執(zhí)行。當(dāng)然一個(gè)SM可以執(zhí)行多個(gè)Wrap预麸。

上圖是SM的組成部分瞪浸，其中綠色的Core是SP，也可以稱為CUDA

core吏祸，CUDA core是一個(gè)單精度微處理器默终，DP Unit是雙精度微處理器。一個(gè)SM里面一般包含32個(gè)DP Unit犁罩，64個(gè)CUDA core（SP）齐蔽。此外SM換包括LD/ST單元，LD/ST主要用于內(nèi)存操作床估，SM中另外要給組件SFU（Special function unit）部分含滴，主要用于執(zhí)行一些編譯器內(nèi)建的特殊函數(shù)，例如cos等丐巫。

SP（Streamingprocessor）：流處理器谈况，GPU中專門用于數(shù)學(xué)（整形，浮點(diǎn)）計(jì)算的處理單元递胧，多個(gè)SP與其他組件一起組成一個(gè)SM碑韵。SP相當(dāng)于一個(gè)微型處理器，基本結(jié)構(gòu)圖如下所示：

GPU內(nèi)存架構(gòu)

GPU內(nèi)存邏輯結(jié)構(gòu)如下：

LocalMemory：本地內(nèi)存华蜒，屬于線程私有內(nèi)存，只能有當(dāng)前線程訪問豁遭，其他線程不可訪問叭喜，當(dāng)線程內(nèi)的寄存器不夠使用時(shí)，會(huì)動(dòng)用該部分內(nèi)存蓖谢。

sharedMemory：共享內(nèi)存捂蕴，可以被同一個(gè)block中的所有線程訪問，不同的線程可以通過共享內(nèi)存進(jìn)行通信闪幽。一般一個(gè)block內(nèi)會(huì)對(duì)應(yīng)一塊共享內(nèi)存启绰，該共享內(nèi)存只能被block內(nèi)的線程訪問（Block中的線程會(huì)被分配給同一個(gè)SM執(zhí)行）。

Global memory：全局內(nèi)存（占據(jù)顯卡內(nèi)存的絕大部分）沟使，CPU和GPU都可以訪問（CPU通過PCIe總線進(jìn)行訪問）委可，所有線程都可以訪問Global memory。

常量內(nèi)存：特殊類型的全局內(nèi)存腊嗡，Grid內(nèi)的所有線程都可以只讀訪問着倾。

紋理內(nèi)存：特殊類型的全局內(nèi)存，Grid內(nèi)的所有線程都可以只讀訪問燕少。

線程組織方式

線程首先被組織為線程塊（thread

block）卡者，多個(gè)線程塊被組織成Grid。

線程塊內(nèi)可以進(jìn)行數(shù)據(jù)的共享客们，以及執(zhí)行過程的同步（一個(gè)線程塊內(nèi)的某個(gè)線程可以被掛起崇决，直到其他線程執(zhí)行到同樣的位置）（執(zhí)行流程的同步）。

在線程塊內(nèi)底挫，每個(gè)線程都有一個(gè)線程ID恒傻，根據(jù)線程ID可以進(jìn)行復(fù)雜尋址。

多個(gè)線程塊組成線程Grid建邓，kernel被映射到Grid進(jìn)行執(zhí)行盈厘，因此一個(gè)Kernel可以使用的線程數(shù)量非常巨大。

同一個(gè)Grid內(nèi)不同Block的線程彼此之間不能通訊和同步官边。在Grid內(nèi)的每個(gè)Block都有要唯一性的ID進(jìn)行標(biāo)識(shí)沸手，根據(jù)Block ID可以進(jìn)行復(fù)雜尋址。

硬件在調(diào)度時(shí)將block劃分為固定數(shù)量的Wrap注簿，同一個(gè)block對(duì)應(yīng)的wrap會(huì)被調(diào)度到某一個(gè)SM執(zhí)行契吉，不同Block中的線程可能被同一個(gè)SM執(zhí)行，也可能被不同的SM執(zhí)行诡渴。

CUDA編碼

CUDA中捐晶，編碼上通過編寫kernel函數(shù)來讓GPU完成計(jì)算任務(wù)。

CUDA函數(shù)

CUDA中的函數(shù)分為三類。

1：只能在CPU側(cè)調(diào)用和執(zhí)行的函數(shù)租悄，該類函數(shù)通過在函數(shù)前添加__host__限定詞限制。

2：在主機(jī)側(cè)被調(diào)用恩袱，在GPU設(shè)備側(cè)被執(zhí)行的函數(shù)泣棋，這類函數(shù)通過在函數(shù)聲明前添加__global__限定詞限制。對(duì)于這類函數(shù)畔塔，返回值只能時(shí)void潭辈，并且不支持遞歸，參數(shù)的大小也被限制在256個(gè)字節(jié)以內(nèi)澈吨，該類函數(shù)內(nèi)也不能聲明靜態(tài)變量把敢。對(duì)于被__global__修飾的函數(shù)，在調(diào)用時(shí)谅辣，需要指明執(zhí)行線程的配置修赞。

3：只能在GPU設(shè)備側(cè)被調(diào)用的函數(shù)，這類函數(shù)通過__device__修飾桑阶。被__device__修飾的函數(shù)無法獲取函數(shù)地址柏副。

CUDA變量

對(duì)于變量而言，通過不同的修飾限定詞蚣录，也區(qū)分為以下幾種

1：被__device__修飾的變量割择，該類變量駐留在GPU設(shè)備全局內(nèi)存空間中，在應(yīng)用的整個(gè)生命周期中有效萎河，Grid內(nèi)的所有線程都可以訪問該變量荔泳，主機(jī)側(cè)通過runtime庫也可以訪問該變量。

2：被__constant__修改的變量虐杯，該類變量駐留在GPU側(cè)的常量內(nèi)存空間中玛歌，在應(yīng)用的整個(gè)聲明周期有效，Grid內(nèi)的所有線程都可以訪問該變量擎椰，主機(jī)側(cè)通過runtime庫也可以訪問該變量沾鳄。

3：被__shared__修飾的變量，該類變量駐留在Block線程塊的共享內(nèi)存中确憨，在Block聲明周期內(nèi)有效译荞，只有同Block內(nèi)的線程可以訪問。

CUDA函數(shù)的執(zhí)行

通過__global__修飾的kernel函數(shù)休弃，在主機(jī)側(cè)被調(diào)用時(shí)吞歼，需要指定該函數(shù)如何在GPU側(cè)被執(zhí)行。

例如對(duì)于如下一個(gè)kernel函數(shù)

__global__ void

Func(float *param)塔猾，在調(diào)用的時(shí)候形式如下：

Func<<<Dg,

Db, Ns, S>>>(param);

其中<<<Dg, Db, Ns, S >>>部分描述了GPU上執(zhí)行的線程組織方式篙骡。

Dg是一個(gè)dim3（x,y,z）的數(shù)據(jù)類型或int型，用于指定Grid的維數(shù)和大小。如果時(shí)int型糯俗，表示是一個(gè)一維的組織結(jié)構(gòu)尿褪。

Db是一個(gè)dim3的數(shù)據(jù)類型或int型，用于指定Block的維數(shù)和大小得湘。如果是int型杖玲，表示是一個(gè)一維的組織結(jié)構(gòu)。

Ns是一個(gè)size_t類型的數(shù)據(jù)淘正，用于指定每個(gè)block中可動(dòng)態(tài)分配的共享內(nèi)存的字節(jié)數(shù)量摆马，默認(rèn)為0.

S：cuda流類型，默認(rèn)為0.

如下的調(diào)用方式

dim3 grid(3,2,1), block(4,3,1);

kernel_name<<<grid,

block>>>(...);

???????? 表示一個(gè)Grid中有3*2*1個(gè)Block鸿吆，在（x,y,z）三個(gè)方向上排列方式分表為3/2/1;一個(gè)Block中有4*3*1個(gè)線程囤采，在(x,y,z)三個(gè)方向上排列方式為4/3/1.

如下的kernel調(diào)用方式：

???????? kernel_name<<<5,8>>>(...);

???????? 表示一個(gè)Grid中有5個(gè)Block，在(x,y,z)三個(gè)方向上的分布方式為5/1/1;在一個(gè)Block中有8個(gè)線程惩淳，在(x,y,z)三個(gè)方向上的分布方式為8/1/1蕉毯。

在kernel函數(shù)中，可以通過內(nèi)置變量計(jì)算線程ID

threadIdx.[x, y, z]表示Block內(nèi)Thread的編號(hào)

blockIdx.[x, y, z]表示Gird內(nèi)Block的編號(hào)

blockDim.[x, y, z]表示Block的維度思犁，也就是Block中每個(gè)方向上的Thread的數(shù)目

gridDim.[x, y, z]表示Gird的維度恕刘，也就是Grid中每個(gè)方向上Block的數(shù)目

對(duì)于kernel_name<<<4, 8>>>(...)這樣一維的結(jié)構(gòu)，如下所示：

線程ID = blockIdx.x *

blockDim.x + threadIdx.x

對(duì)于如下二維結(jié)構(gòu)

dim grid(4,1,1),

block(2,2,1);

kernel_name<<<grid,

block>>>(...)

線程ID計(jì)算如下：

int blockId =

blockIdx.x + blockId.y * gridDim.x;

int threadId =

blockId * (blockDim.x * blockDim.y) + (threadIdx.y *blockDim.x) + threadIdx.x;

對(duì)于如下多維結(jié)構(gòu)kernel函數(shù)抒倚，線程ID計(jì)算方式如下：

int blockId =

blockIdx.x + blockIdx.y * gridDim.x + gridDim.x * gridDim.y * blockIdx.z;

int threadIc =

blockId * (blockDim.x * blockDim.y * blockDim.z)

?????????????????????? + (threadIdx.z *(blockDim.x * blockDim.y))

?????????????????????? + (threadIdx.y *blockDim.x) + threadIdx.x;

內(nèi)置變量

在編寫CUDA代碼時(shí)褐着，有一些內(nèi)置變量可以直接使用。內(nèi)置變量不允許獲取變量地址托呕，也不允許被賦值含蓉。

CUDA特殊數(shù)據(jù)類型

Dim3數(shù)據(jù)類型是CUDA中比較特殊的數(shù)據(jù)類型，Dim3是基于uint3類型的3維結(jié)構(gòu)體项郊，定義一個(gè)dim3類型馅扣，默認(rèn)取值都為1。

CUDA運(yùn)行庫函數(shù)

CUDA本身提供了一些可以使用的函數(shù)着降，主要有以下幾類

1：設(shè)備管理類的函數(shù)差油，例如cudaGetDeviceCount等，這類函數(shù)在主機(jī)側(cè)被調(diào)用任洞，用于查詢?cè)O(shè)備相關(guān)的信息蓄喇。

2：內(nèi)存管理相關(guān)函數(shù)，例如cudaMalloc等函數(shù)交掏，這類函數(shù)允許用戶申請(qǐng)?jiān)O(shè)備內(nèi)存以及在主機(jī)內(nèi)存于設(shè)備內(nèi)存之間進(jìn)行數(shù)據(jù)傳遞妆偏。

3：流管理相關(guān)函數(shù)。用于創(chuàng)建和銷毀流盅弛。

4：事件管理相關(guān)函數(shù)钱骂，事件的創(chuàng)建和銷毀叔锐，跟蹤等。

5：原子函數(shù)等其他函數(shù)见秽。

CUDA流與事件

????一個(gè)典型的CUDA程序愉烙，應(yīng)該是GPU完成一部分工作，CPU完成一部分工作解取，當(dāng)CPU把任務(wù)交給GPU執(zhí)行時(shí)步责，有兩種策略，第一種策略是等待GPU完成后肮蛹，CPU繼續(xù)執(zhí)行其他任務(wù)（同步機(jī)制）勺择；另外一種策略是创南，CPU把任務(wù)交給GPU后伦忠，不做等待，繼續(xù)執(zhí)行需要在CPU側(cè)完成的任務(wù)稿辙，經(jīng)過一定時(shí)間后昆码，通過查詢的方式查詢GPU是否已經(jīng)完成了對(duì)應(yīng)的任務(wù)，或者當(dāng)GPU完成任務(wù)后邻储，通知CPU任務(wù)已經(jīng)完成赋咽，為了支持第二種策略，CUDA提供給了流和事件機(jī)制吨娜。

CUDA流

CUDA流可以看作是在GPU上執(zhí)行任務(wù)的一個(gè)隊(duì)列脓匿，并且該隊(duì)列中的操作任務(wù)可以按照指定的順序執(zhí)行，并且可以創(chuàng)建多個(gè)流（多個(gè)操作隊(duì)列）宦赠，在流和流之間也是可以并行執(zhí)行陪毡。

在CUDA中，流被分為兩種類型勾扭，隱式流（或者說是匿名流）毡琉，該流沒有名稱，無法直接進(jìn)行控制妙色，在執(zhí)行上默認(rèn)與CPU之間是同步模式桅滋；顯式流（需要主動(dòng)的聲明和創(chuàng)建對(duì)應(yīng)的流），可以對(duì)該流進(jìn)行代碼層的直接控制（可以創(chuàng)建同步流身辨，CPU需要等待GPU的執(zhí)行結(jié)果丐谋，也可以創(chuàng)建異步流，CPU不需要等待GPU執(zhí)行結(jié)果）煌珊。

不同流中的操作可以并行執(zhí)行笋鄙，但同一個(gè)流中的操作無法并行執(zhí)行。Hyper-Q技術(shù)的出現(xiàn)可以支持多個(gè)流實(shí)現(xiàn)完全的并行執(zhí)行怪瓶，邏輯圖如下：

流也可以定義不同的優(yōu)先級(jí)萧落，從而讓高優(yōu)先級(jí)的流可以得到更多的執(zhí)行機(jī)會(huì)践美。

在CUDA中，主要通過以下幾個(gè)接口完成流的創(chuàng)建和使用

1：創(chuàng)建一個(gè)流cudaError_t cudaStreamCreate(cudaStream_t* stream);

2：查詢?cè)摿鲗?duì)應(yīng)的操作是否已經(jīng)完成cudaError_t cudaStreamQuery(cudaStream_t stream)

3：cudaError_t cudaStreamSyncronize(cudaStream_t stream);阻塞當(dāng)前進(jìn)程執(zhí)行找岖，直到設(shè)備上對(duì)應(yīng)的流中的所有操作都完成陨倡。

4：cudaError_t cudaStreamDestroy(cudaStream_t stream); 釋放流所占用的資源。

CUDA事件

CUDA中事件有兩個(gè)作用许布。

1：可以在流的執(zhí)行過程中添加標(biāo)記點(diǎn)兴革，檢查執(zhí)行流是否已經(jīng)到達(dá)對(duì)應(yīng)點(diǎn)，可用于執(zhí)行過程的等待和測(cè)試蜜唾，作用類似于cudaStreamSyncronize杂曲。

2：可以插入到不同的流中，在不同的流中進(jìn)行操作過程的同步袁余，也即可以通過主機(jī)端對(duì)設(shè)備端執(zhí)行過程進(jìn)行操控擎勘。

3：可用于執(zhí)行過程的執(zhí)行時(shí)長統(tǒng)計(jì)。

主要有以下幾個(gè)接口可以使用

1：cudaError_t

cudaEventCreate(cudaEvent_t* event);創(chuàng)建要給事件并得到一個(gè)句柄颖榜。

2：cudaError_t

cudaEventRecord(cudaEvent_t event, CUstream stream);記錄一個(gè)事件棚饵，當(dāng)給定了流參數(shù)后，當(dāng)流中所有操作都完成后掩完，該事件被記錄噪漾。

3：cudaError_t

cudaEventQuery(cudaEvent_t event);查詢事件是否已經(jīng)被記錄。

4：cudaError_t cudaEventSyncronize(cudaEvent_t

event);阻塞執(zhí)行過程且蓬，直到對(duì)應(yīng)事件已經(jīng)發(fā)生欣硼。cudaError_t cudaEventDestroy(cudaEvent_t event);銷毀事件占用的資源。

CUDA代碼編譯

編寫的CUDA程序代碼恶阴，不能直接使用已有的編譯器進(jìn)行編譯诈胜，需要使用nvcc工具先進(jìn)行處理，從編寫的代碼中分離出設(shè)備側(cè)代碼存淫，并將設(shè)備側(cè)代碼編譯為二進(jìn)制格式耘斩。

分離出的代碼文件可以直接使用標(biāo)準(zhǔn)編譯器進(jìn)行編譯，然后于CUDA runtime庫鏈接生成可執(zhí)行程序桅咆。

CUDA代碼樣例

匿名流樣例

一個(gè)典型的CUDA程序如下：

1：定義可以從主機(jī)側(cè)調(diào)用的kernel函數(shù)

__global__宏限制addKernel函數(shù)在主機(jī)側(cè)調(diào)用括授，在GPU側(cè)被執(zhí)行。傳遞的參數(shù)a和b做為加法運(yùn)算的“加數(shù)”和“被加數(shù)”岩饼，參數(shù)c做為結(jié)果荚虚。內(nèi)部計(jì)算時(shí)，通過threadIdx獲取執(zhí)行addKernel函數(shù)的線程籍茧，讓每個(gè)線程計(jì)算對(duì)應(yīng)數(shù)組中的某一個(gè)版述。

2：通過CUDA運(yùn)行庫函數(shù)，獲取和設(shè)置用于計(jì)算的GPU設(shè)備

3：調(diào)用cuda相關(guān)庫函數(shù)在GPU上分配用于計(jì)算的內(nèi)存空間

4：調(diào)用cuda接口函數(shù)將主機(jī)內(nèi)存中的數(shù)據(jù)拷貝到device GPU側(cè)內(nèi)存中

5：調(diào)用kernel函數(shù)寞冯，在GPU上進(jìn)行計(jì)算

6：調(diào)用CUDA接口等待GPU設(shè)備側(cè)完成計(jì)算渴析，或執(zhí)行其主機(jī)側(cè)代碼晚伙。該函數(shù)會(huì)阻塞，直到GPU側(cè)完成計(jì)算后才返回俭茧。

7：調(diào)用CUDA接口將GPU側(cè)執(zhí)行的結(jié)果拷貝會(huì)主機(jī)內(nèi)存

8：釋放分配的GPU側(cè)內(nèi)存咆疗，釋放占用的GPU設(shè)備

多流樣例

?著作權(quán)歸作者所有,轉(zhuǎn)載或內(nèi)容合作請(qǐng)聯(lián)系作者

人面猴
序言：七十年代末，一起剝皮案震驚了整個(gè)濱河市母债，隨后出現(xiàn)的幾起案子午磁，更是在濱河造成了極大的恐慌，老刑警劉巖毡们，帶你破解...
沈念sama閱讀 216,470評(píng)論 6贊 501
死咒
序言：濱河連續(xù)發(fā)生了三起死亡事件迅皇，死亡現(xiàn)場(chǎng)離奇詭異，居然都是意外死亡衙熔，警方通過查閱死者的電腦和手機(jī)登颓，發(fā)現(xiàn)死者居然都...
沈念sama閱讀 92,393評(píng)論 3贊 392
救了他兩次的神仙讓他今天三更去死
文/潘曉璐我一進(jìn)店門，熙熙樓的掌柜王于貴愁眉苦臉地迎上來青责，“玉大人挺据，你說我怎么就攤上這事取具〔绷ィ” “怎么了？”我有些...
開封第一講書人閱讀 162,577評(píng)論 0贊 353
道士緝兇錄：失蹤的賣姜人
文/不壞的土叔我叫張陵暇检，是天一觀的道長产阱。經(jīng)常有香客問我，道長块仆，這世上最難降的妖魔是什么构蹬？我笑而不...
開封第一講書人閱讀 58,176評(píng)論 1贊 292
?港島之戀（遺憾婚禮）
正文為了忘掉前任，我火速辦了婚禮悔据，結(jié)果婚禮上庄敛，老公的妹妹穿的比我還像新娘。我一直安慰自己科汗，他們只是感情好藻烤，可當(dāng)我...
茶點(diǎn)故事閱讀 67,189評(píng)論 6贊 388
惡毒庶女頂嫁案：這布局不是一般人想出來的
文/花漫我一把揭開白布。她就那樣靜靜地躺著头滔，像睡著了一般怖亭。火紅的嫁衣襯著肌膚如雪。梳的紋絲不亂的頭發(fā)上坤检，一...
開封第一講書人閱讀 51,155評(píng)論 1贊 299
城市分裂傳說
那天兴猩，我揣著相機(jī)與錄音，去河邊找鬼早歇。笑死倾芝，一個(gè)胖子當(dāng)著我的面吹牛讨勤，可吹牛的內(nèi)容都是我干的。我是一名探鬼主播晨另，決...
沈念sama閱讀 40,041評(píng)論 3贊 418
雙鴛鴦連環(huán)套：你想象不到人心有多黑
文/蒼蘭香墨我猛地睜開眼悬襟，長吁一口氣：“原來是場(chǎng)噩夢(mèng)啊……” “哼！你這毒婦竟也來了拯刁？” 一聲冷哼從身側(cè)響起脊岳，我...
開封第一講書人閱讀 38,903評(píng)論 0贊 274
萬榮殺人案實(shí)錄
序言：老撾萬榮一對(duì)情侶失蹤，失蹤者是張志新（化名）和其女友劉穎垛玻，沒想到半個(gè)月后割捅，有當(dāng)?shù)厝嗽跇淞掷锇l(fā)現(xiàn)了一具尸體，經(jīng)...
沈念sama閱讀 45,319評(píng)論 1贊 310
?護(hù)林員之死
正文獨(dú)居荒郊野嶺守林人離奇死亡帚桩，尸身上長有42處帶血的膿包…… 初始之章·張勛以下內(nèi)容為張勛視角年9月15日...
茶點(diǎn)故事閱讀 37,539評(píng)論 2贊 332
?白月光啟示錄
正文我和宋清朗相戀三年亿驾，在試婚紗的時(shí)候發(fā)現(xiàn)自己被綠了。大學(xué)時(shí)的朋友給我發(fā)了我未婚夫和他白月光在一起吃飯的照片账嚎。...
茶點(diǎn)故事閱讀 39,703評(píng)論 1贊 348
活死人
序言：一個(gè)原本活蹦亂跳的男人離奇死亡莫瞬，死狀恐怖，靈堂內(nèi)的尸體忽然破棺而出郭蕉，到底是詐尸還是另有隱情疼邀，我是刑警寧澤，帶...
沈念sama閱讀 35,417評(píng)論 5贊 343
?日本核電站爆炸內(nèi)幕
正文年R本政府宣布召锈，位于F島的核電站旁振，受9級(jí)特大地震影響，放射性物質(zhì)發(fā)生泄漏涨岁。R本人自食惡果不足惜拐袜，卻給世界環(huán)境...
茶點(diǎn)故事閱讀 41,013評(píng)論 3贊 325
男人毒藥：我在死后第九天來索命
文/蒙蒙一、第九天我趴在偏房一處隱蔽的房頂上張望梢薪。院中可真熱鬧蹬铺，春花似錦、人聲如沸秉撇。這莊子的主人今日做“春日...
開封第一講書人閱讀 31,664評(píng)論 0贊 22
一樁弒父案，背后竟有這般陰謀
文/蒼蘭香墨我抬頭看了看天上的太陽畜疾。三九已至赴邻，卻和暖如春，著一層夾襖步出監(jiān)牢的瞬間啡捶，已是汗流浹背姥敛。一陣腳步聲響...
開封第一講書人閱讀 32,818評(píng)論 1贊 269
情欲美人皮
我被黑心中介騙來泰國打工，沒想到剛下飛機(jī)就差點(diǎn)兒被人妖公主榨干…… 1. 我叫王不留瞎暑，地道東北人彤敛。一個(gè)月前我還...
沈念sama閱讀 47,711評(píng)論 2贊 368
代替公主和親
正文我出身青樓与帆，卻偏偏與公主長得像，于是被迫代替她去往敵國和親墨榄。傳聞我的和親對(duì)象是個(gè)殘疾皇子玄糟，可洞房花燭夜當(dāng)晚...
茶點(diǎn)故事閱讀 44,601評(píng)論 2贊 353