CUDA Materials Collection

Other than NVIDIA CUDA Programming Guide
NVIDIA CUDA Basics [PDF]
NVIDIA CUDA Basics [WMV]
NVIDIA Optimizing CUDA [PDF]
NVIDIA Advanced CUDA Webinar: Memory Optimizations [WMV]

NVIDIA: GPU Computing Online Seminar

Threads, Blocks and Grids

From NVIDIA Optimizing CUDA

“Experiment” mode discovers and saves optimal configuration

Blocks per Grid Heuristics

• Number of blocks should be more than number of multiprocessors, so all multiprocessors have at least one block to execute.
  
• Number of blocks divided by number of multiprocessors should be more than 2
  
  • Multiple blocks can run concurrently in a multiprocessor
    
  • Blocks that aren’t waiting at a __syncthreads() keep the hardware busy
    
  • Subject to resource availability –registers, shared memory
    
• Number of blocks that is more than 100 will scale to future devices

Optimizing threads per block

• Choose threads per block as a multiple of warp size to avoid wasting computation on under-populated warps and to facilitate coalescing
  
• Run as many warps as possible per multiprocessor (hide latency)
  
• Multiprocessor can run up to 8 blocks at a time
  
• Heuristics
  
  • Minimum: 64 threads per block (Only if multiple concurrent blocks)
    
  • 192 or 256 threads a better choice (Usually still enough regs to compile and invoke successfully)
    
  • However, this all depends on your computation, so experiment!
    

Zero Copy (Mapped Memory): Directly Access Host Memory from The Device

Came across with this Zero Copy which has just been introduced in CUDA 2.2 when I was looking for the concept of threads and blocks.

Quoted from NVIDIA Optimizing CUDA

• Access host memory directly from device code
• Transfers implicitly performed as needed by device code
• All set-up is done on host using mapped memory

What should be considered when using Zero Copy

• Zero copy will always be a win for integrated devices that utilize CPU memory (check this using the integratedfield in cudaDeviceProp)
  
• Zero copy will be faster if data is only read/written from/to global memory once: Copy input to GPU, One kernel run, Copy output to CPU
  
• Potentially easier and faster alternative to using cudaMemcpyAsync
  
• Current devices use pointers that are 32-bit so there is a limit of 4GB per context

Quoted from section 3.2.5.3 Mapped Memory in CUDA Programming Guide 2.2

On some devices, a block of page-locked host memory can also be mapped into the device’s address space by passing flag cudaHostAllocMapped to cudaHostAlloc(). Such a block has therefore two addresses: one in host memory and one in device memory. The host memory pointer is returned by cudaHostAlloc() and the device memory pointer can be retrieved using cudaHostGetDevicePointer() and used to access the block from within a
kernel.

Accessing host memory directly from within a kernel has several advantages:

• There is no need to allocate a block in device memory and copy data between this block and the block in host memory; data transfers are implicitly performed as needed by the kernel;
• There is no need to use streams (see Section 3.2.6.1) to overlap data transfers with kernel execution; the kernel-originated data transfers automatically overlap with kernel execution.
  

Since mapped page-locked memory is shared between host and device however, the application must synchronize memory accesses using streams or events (see Section 3.2.6) to avoid any potential read-after-write, write-after-read, or write-after-
write hazards.

A block of page-locked host memory can be allocated as both mapped and portable (see Section 3.2.5.1), in which case each host thread that needs to map the block to its device address space must call cudaHostGetDevicePointer() to retrieve a device pointer, as device pointers will generally differ from one host thread to the other.

To be able to retrieve the device pointer to any mapped page-locked memory within a given host thread, page-locked memory mapping must be enabled by calling cudaSetDeviceFlags() with the cudaDeviceMapHost flag before any other CUDA operations is performed by the thread. Otherwise, cudaHostGetDevicePointer() will return an error.

cudaHostGetDevicePointer() also returns an error if the device does not support mapped page-locked host memory.

Applications may query whether a device supports mapped page-locked host memory or not by calling cudaGetDeviceProperties() and checking the canMapHostMemory property.

CUDA Wrapping

Quoted from Section 4.1 of CUDA Programming Guide 2.2.1: A Set of SIMT Multiprocessors with On-Chip Shared Memory

The CUDA architecture is built around a scalable array of multithreaded Streaming Multiprocessors (SMs). When a CUDA program on the host CPU invokes a kernel grid, the blocks of the grid are enumerated and distributed to multiprocessors with available execution capacity as illustrated in Figure 4-1. The threads of a thread block execute concurrently on one multiprocessor. As thread blocks terminate, new blocks are launched on the vacated multiprocessors.

A multiprocessor consists of eight Scalar Processor (SP) cores, two special function units for transcendentals, a multithreaded instruction unit, and on-chip shared memory. The multiprocessor creates, manages, and executes concurrent threads in hardware with zero scheduling overhead. It implements the __syncthreads() barrier synchronization intrinsic with a single instruction. Fast barrier synchronization together with lightweight thread creation and zero-overhead thread scheduling efficiently support very fine-grained parallelism, allowing, for example, a low granularity decomposition of problems by assigning one thread to each data element (such as a pixel in an image, a voxel in a volume, a cell in a grid-based computation).

To manage hundreds of threads running several different programs, the multiprocessor employs a new architecture we call SIMT (single-instruction, multiple-thread). The multiprocessor maps each thread to one scalar processor core, and each scalar thread executes independently with its own instruction address and register state. The multiprocessor SIMT unit creates, manages, schedules, and executes threads in groups of 32 parallel threads called warps. (This term originates from weaving, the first parallel thread technology. A half-warp is either the first or second half of a warp.) Individual threads composing a SIMT warp start together at the same program address but are otherwise free to branch and execute independently.

When a multiprocessor is given one or more thread blocks to execute, it splits them into warps that get scheduled by the SIMT unit. The way a block is split into warps is always the same; each warp contains threads of consecutive, increasing thread IDs with the first warp containing thread 0. Section 2.2 describes how thread IDs relate to thread indices in the block.

Every instruction issue time, the SIMT unit selects a warp that is ready to execute and issues the next instruction to the active threads of the warp. A warp executes one common instruction at a time, so full efficiency is realized when all 32 threads of a warp agree on their execution path. If threads of a warp diverge via a data-dependent conditional branch, the warp serially executes each branch path taken, disabling threads that are not on that path, and when all paths complete, the threads converge back to the same execution path. Branch divergence occurs only within a warp; different warps execute independently regardless of whether they are executing common or disjointed code paths.

SIMT architecture is akin to SIMD (Single Instruction, Multiple Data) vector organizations in that a single instruction controls multiple processing elements. A key difference is that SIMD vector organizations expose the SIMD width to the software, whereas SIMT instructions specify the execution and branching behavior of a single thread. In contrast with SIMD vector machines, SIMT enables programmers to write thread-level parallel code for independent, scalar threads, as well as data-parallel code for coordinated threads. For the purposes of correctness, the programmer can essentially ignore the SIMT behavior; however, substantial performance improvements can be realized by taking care that the code seldom requires threads in a warp to diverge. In practice, this is analogous to the role of cache lines in traditional code: Cache line size can be safely ignored when designing for correctness but must be considered in the code structure when designing for peak performance. Vector architectures, on the other hand, require the software to coalesce loads into vectors and manage divergence manually.

Not quite understood but was trying to read more.

Out to Swim (Again)

I intended to go swimming this afternoon... Really intended!

เมื่อตอนกลางวัน เราตั้งใจจะไปว่ายน้ำ ตั้งใจมากๆ

I left my house at nearly four, felt good that the sunlight was not so strong like the other days. I was not sure how long I had been swimming. Perhaps about half an hour. Then the rain came. Though it was an indoor swimming pool, the staff called everyone to get out of the pool. Do not ask me why. I do not know.

เราออกจากบ้านเมื่อตอนเกือบสี่โมง ดีที่วันนี้แดดไม่จัดมาก ไปถึงก็ว่ายน้ำ ไม่แน่ใจว่าได้ว่ายอยู่นานเท่าไหร่ ประมาณครึ่งชั่วโมงได้มั้ง แล้วฝนก็เทลงมา จริงๆ สระมันอยู่ในตึก แต่คนดูแลก็เรียกทุกคนขึ้นจากสระ อย่าถามว่าทำไม หนูไม่ยู้

At least I got some exercises after all. I walked to the pool, swam, and walked back. hehe

เอาน่ะ อย่างน้อยหนูก็ได้ออกกำลังกายแหละ เริ่มจากเดินไปสระ ว่ายน้ำ แล้วก็เดินกลับ อิอิ

Summer Class

It has been three days since the class began. I have class everyday this quarter and there are only six weeks to the end of quarter. The weather is so hot that I do not want to step outside. :(

These days I have been addicted to comics that Pao (my senior) has introduced me. It is bad. I should have been focusing on studying and my research.

At McDonald's

Today, I had lunch at McDonald's. It was the first meal that I ate out since I got back to Ruston. It was because I wanted somewhere to hide from the sun. I had and appointment with the landlady of my new apartment at 10 AM. So I left my house around 9:10, walked, and got there around 9:40. It was so hot because of the sunlight. After having a talk with the landlady I walked back. The landlady actually offered me a ride but I would like to stop by somewhere (which i realized that it was closed on Wednesdays afterward :( ).

วันนี้ไปกินกลางวันที่แมคโดนัลด์มา เป็นมื้อแรกที่ออกไปกินข้างนอกหลังจากกลับมาถึง Ruston ทั้งนี้ทั้งนั้น เพราะไอ้เอ๋อยากหาที่หลบแดด เดินไปอพาร์ตเมนต์ใหม่ตั้งแต่เก้าโมงสิบ เพราะนัดกับเจ้าของอพาร์ตเมนต์ไว้ตอนสิบโมง ไปถึงนั่นเก้าโมงสี่สิบ อากาศโคตรร้อน แดดแรงมาก พอคุยกับเจ้าของเสร็จแล้วก็เดินกลับ จริงๆ เจ้าของบ้านก็บอกว่าจะมาส่ง แต่ไอ้เอ๋จะแวะซื้อของก่อน (ซึ่งมารู้ที่หลังว่ามันปิดวันพุธ :( เซ็งเป็ด! )

I stopped by County Market around 10:30, grabbed some bread, instant noodles, etc. Then I crossed the street, walked into McDonald's and made an order, paid for it and waited. I waited in front of the counter for couple of minutes with no one got me my meal. :(

ทีนี้ก็เลยแวะ County Market ตอนประมาณสิบโมงครึ่ง ซึ้อขนมปังกับบะหมี่กึ่งสำเร็จรูป แล้วก็อย่างอื่นอีกหน่อย แล้วก็เลยข้ามถนนไปแมค สั่งๆ จ่ายๆ แล้วยืนรอ รอเงกอยู่หน้าเคาน์เตอร์ตั้งนานก็ยังไม่มีใครเอาแฮมเบอร์เกอร์ตูมาให้

However, there at last was a black girl realized that I was waiting for something. She asked me and I told her I was waiting for my burger and fries. She got me my cheeseburger, and another black girl told me to wait for the fries. The were not done yet. After couple of minutes, and the fries were getting done, the first black girl asked me again because she saw me still waiting there.

และแล้วก็มีสาวผิวหมึกคนหนึ่งเห็นไอ้เอ๋รออะไรอยู่ พี่แกเลยเข้ามาถาม ไอ้เอ๋ก็เลยบอกไปว่ารอเบอร์เกอร์กับเฟรนช์ฟรายของตูอยู่ พี่แกก็หยิบชีสเบอร์เกอร์มาให้ สาวหมึกอีกคนก็บอกให้รอเฟรนช์ฟรายก่อน กำลังทอดอยู่ ไม่กี่นาทีต่อมา เฟรนช์ฟรายเสร็จแล้ว สาวผิวหมึกคนแรกก็มาถามต่อ ว่าตูรออะไรอีก

"My fries..." I told her.

"เฟรนช์ฟรายหนูล่ะ..."

I finally got my fries with bunch of ketchup packets. :)

ไอ้เอ๋ได้เฟรนช์ฟรายในที่สุด พร้อมกับห่อซอสมะเขืออีกกำนึง :)

Let's Jump into the Water!

I just got back from the swimming pool. It was good to get some exercises but I was so tired. Tech's swimming pool is ok, not bad but not too good. At least, it is bigger than the ones in my village back in Thailand.

เพิ่งไปว่ายน้ำมาแหละ ได้ออกกำลังบ้างก็ดีเหมือนกัน แต่โคตรเหนื่อยเลย สระว่ายน้ำที่นี่ก็ใช้ได้ ไม่ได้แย่ แต่ก็ไม่ได้ดีมาก อย่างน้อยมันก็ใหญ่กว่าสระว่ายน้ำในหมู่บ้านที่เมืองไทยละกัน

My sister had told me to go to swim sometimes. She said that, for me, it is better than other sports since I have varicose vein (blood circulation problem) on my legs. It used to cause me pain when I played badminton or other sports. And I also have allergy. Let us see if I can get rid of it!

โอ๋เคยบอกว่าให้ไปว่ายน้ำบ้าง เพราะเรามีเส้นเลือดขอดที่ขา ทำให้ปวดเวลาเล่นกีฬาอย่างอื่น แล้วก็ยังเป็นภูมิแพ้อีก ก็ดูแล้วกันว่าจะหายมั้ย

Small Party at Gan't Apartment

There was a small party (only Thais) at Gan's apartment last night. There was nothing but hamburgers and salad. Aig and Mon had prepared the hamburgers for us. They were really nice, wern't they? :) There was some alcoholic drink as well. Mon and I did not touch it for sure.

It was my first time to grill hamburgers as well. (Actually Aid, Mon and Gan did it.) They sprayed some oil, put some salt, black pepper and thyme on the beef loafs. I had no idea that there would be thyme on the hamburgers. Because there was rain all day, it was hard to lit the stove. But Gan did it at last. The hamburgers were good. That might be because we did it ourselves. :)

It's SUMMER!

กลับมาได้สามวัน เจออากาศร้อนเหยียบร้อยกันทุกวัน แทบกรี๊ด

ตอนนี้ที่ Ruston อากาศร้อนมาก เปิดแอร์กันทั้งวันทั้งคืน ค่าไฟบาน แล้วนี่มันคงจะร้อนกันไปอีกอย่างน้อยสามเดือน ละลายดีกว่า

4 Weeks in BKK

Four weeks in Bangkok were such a short period. I did many things and missed many things.

I have learned something about genetics with Biotec folk, worked with them for many days. I did not only have Thai food but Vietnamese, Japanese, Chinese, especially seafood which I had really missed.

I did not go travel, did not have enough time to go and no one had free time to go with me. But that was not a big deal. My purpose was to see my family.

New Interest of My Research

I just talked to the staffs of Biotec in Bangkok. They are interested in exploiting GPGPU in their application. They would like to apply some algorithms (now they are using EM) to infer haploid genotype. Anyhow this genetic area is new to me. I will study about it during summer. Hopefully I will be able to gain something to write about before fall.

Note:
http://en.wikipedia.org/wiki/Haplotype
http://en.wikipedia.org/wiki/EM_algorithm

keywords: monte carlo, EM algo (Expectation-maximization algo), expectation maximization haplotyping

BKK 2... Reunion

My cousins, my sister and I got together yesterday. We had Sukiyaki for lunch, played bowling, went shopping and got home to get our grandma and had dinner at a Chinese restaurant.

My niece (almost 3) and my nephew (around 4) has much grown up. They got taller, heavier and more stubborn. Though they were so cute. :)

BKK 1... HOME!

It was a longgggg trip.
I reached home last night. My sister and my cousin were in the airport to pick up another cousin and me. My cousin came from Singapore.

Bangkok has changed a bit as I saw when I was on the way home. At lease there was some huge building I have never seen.

Nothing much for the first day. I was so tired because I didn't sleep for 2 days. I could not sleep on the flight. Do not know why.

Day 426... Ready for Going Home

Today is my last day in Ruston. Tomorrow in early morning (say tonight) I will leave for my home. Yeah!

Home Sweet Home - Mötley Crüe

My First Paper

I just submitted my first paper a minute ago. :D

The title is "Stream Experiments: Toward Latency Hiding in GPGPU", submitted to HiPC 2009.

Alright, now I am ready to go home. Soooo excited!

Day 422... Two Parties

There was a party at Wei and Lee's apartment last Friday. P Noon and I got there around 6 PM. P Noon had prepared some fried eggs with Thai style source. There were couple of egg menus on the table. Moreover, there were some chicken, vegetables, etc.

After eating we played games and talked. The game that I played was a sort of gambling. (Just for fun, no money.) Let's see the pictures



And yesterday, I was in another party since James had invited me to explore Louisiana. It was James' grandmother's birthday party. The place (Sea and Sirloin) was in the middle of nowhere. We got lost because James took the wrong turn. Instead of heading to the south, we went to the north. He realized it when we almost hit I-20 which could lead us back to Ruston in 20 minutes. Then he turned back. We left home around 3:45 and reached the restaurant around 5:20. It was fun though since he talked all the way and told me many things about country lives, red necks, nature of LA and the U.S.

James told me that there were three stages in the party like this, (just for his family) eat, talk and say goodbye at the parking lot. We exactly followed those stages. It reminded me the birthday parties of my grandparents back in Bangkok. They are almost the same but the food and languages. His family was so nice.



From that day, I have realized that Ruston is not too small. It is just small but it is not in the middle of nowhere. Some places, like Friendship, Lucky and Campti (the names of cities) are much smaller. They are so small that I wonder how the people there can live.

Day 419... Let's Party!

It is the last day of classes!!... Wooohooo~

I am looking forward to going back home. But before that, I am going to join some parties here. The Chinese students are going to have a party tomorrow night. And they have invited P Noon and me. It would be fun.

I am going to join James' grandmother's birthday on Saturday as well. He has invited me since I have never had an activity with his family while the others (Aig, Monn and P Noon) have gone on adventures with them last winter. It will be a formal dinner party. I am wondering how it is going to be.

Day 405... I Did It!

Yesterday was one of my great day because I got an email said that I have passed the CS Qualifying Exam... WooHoo! However, I still have to pass the Math one. It would be a hard part as well. :(

I have been very busy in these couple of weeks. The paper submission ends next week, and I have to get the term project done by next week as well, and another term project dues next two weeks, besides the homework and finals... Hell! :(

Day 402... Thai Noodles at Wesley

David Glasspool had invited P Noon to cook Thai food in Wesley last week. We did it today. :)

เดวิด แกลสพูล ชวนพี่นูนไปทำอาหารไทยที่เวสลีเมื่ออาทิตย์ที่แล้ว ก็เลยได้ทำกันวันนี้

It was Pud See Euw. I only helped him chop vegetables. In the meantime, Kevin, Gray and the other guys were cooking as well. I just helped them watch and taste the food. :)

มันคือผัดซีอิ้ว ไอ้เอ๋ช่วยหั่นผักอย่างเดียว ตอนนั้นเควิน เกรย์ กับอีกหลายหนุ่มก็กำลังทำกับข้าวด้วย ไอ้เอ๋ช่วยดูกับช่วยชิม :)

After dinner, we hung out there for a while, talked to people there. Then Nathan ride P Noon home.

หลังกินข้าว ก็นั่งคุยอีกพักนึง แล้วเนธันก็ขับรถพาพี่นูนกลับบ้าน