Houdini OpenCL 筆記(一)


//很簡單的一個案例,獲取idx,根據idx得到P, 運算得到臨時變量pos,把pos寫入P

1  int idx = get_global_id(0);
2  if (idx >= P_length)   
3         return;
4 float3 pos = vload3(idx, P);
5 pos.y += amplitude * sin( length(pos) / period + phase );
6 vstore3(pos, idx, P);

可以得到漣漪的效果

// Houdini 16.5 Masterclass - OpenCL vs VEX 中的一個案例,實現對Volume 進行blur

 1 kernel void kernelName( 
 2                  int height_stride_x, 
 3                  int height_stride_y, 
 4                  int height_stride_z, 
 5                  int height_stride_offset, 
 6                  int height_res_x, 
 7                  int height_res_y, 
 8                  int height_res_z, 
 9                  global float * height,
10                  global float * __scratch
11 #ifdef HAS_mask
12                  , global float * mask
13 #endif
14 )
15 {
16     int gidx = get_global_id(0);
17     int gidy = get_global_id(1);
18     int gidz = get_global_id(2);
19     int idx = height_stride_offset + height_stride_x * gidx
20                                + height_stride_y * gidy
21                                + height_stride_z * gidz;
22 
23 #ifdef HAS_mask
24      if (mask[idx] < 0.5)
25      {
26         __scratch[idx] = height[idx];
27         return;
28      }
29 #endif
30      float total = 0;
31      for (int dx = -1; dx <= 1; dx++)
32      {
33         for (int dy = -1; dy <= 1; dy++)
34         {
35             for (int dz = -1; dz <= 1; dz++)
36             {
37                 int srcidx = height_stride_offset
38                                 + height_stride_x * clamp(gidx+dx, 0, height_res_x-1)
39                                 + height_stride_y * clamp(gidy+dy, 0, height_res_y-1)
40                                 + height_stride_z * clamp(gidz+dz, 0, height_res_z-1);
41                 float src_height = height[srcidx];
42                 
43                 total += src_height;                   
44             }
45         }
46      }
47  
48      __scratch[idx] = total/27;
49 }
50 
51 kernel void writeBack( 
52                  int height_stride_x, 
53                  int height_stride_y, 
54                  int height_stride_z, 
55                  int height_stride_offset, 
56                  int height_res_x, 
57                  int height_res_y, 
58                  int height_res_z, 
59                  global float * height,
60                  global float * __scratch
61 #ifdef HAS_mask
62                  , global float * mask
63 #endif
64 )
65 {
66     int gidx = get_global_id(0);
67     int gidy = get_global_id(1);
68     int gidz = get_global_id(2);
69     int idx = height_stride_offset + height_stride_x * gidx
70                                + height_stride_y * gidy
71                                + height_stride_z * gidz;
72     height[idx] = __scratch[idx];
73 }
View Code

 

 

 

 

 

 

 

height 參數設置時 ,取消勾選了Force Alignment(Force Alignment 是為了簡化代碼,當有多個Volume的時候,認為Resolution和Transform相同),這樣在kernal代碼中,多了height_stride和height_res為前綴的參數。

__scratch 是對height場的拷貝,勾選Force Alignment是因為,Resolution、Voxel Size  以及Transform什么都一樣, height場的idx , 就是我們要對__scratch場操作的idx 。

 mask 勾選Force Alignment的原因也同上,完全可以根據height場的idx,得到mask場的idx。需要注意的是,mask勾選了Optional參數,這樣代碼中,在  #ifdef HAS_mask 和 #endif 中間的語句,只有在有mask場的時候才會執行,否則跳過

// OpenCL對數據的格式要求比較嚴格

比如像這樣的語句:  float d =  abs( x0 ) + abs(y0)  + abs(z0) - 1 。 前面的語句中x0,y0,z0都是float型,按理說 聲明了float 型d,d得到的結果應該是float型,但是由於等式右邊 1 是整形,最終導致d得到的結果是整形。這個在16.5中存在這樣的問題,不知道以后會不會修復。

上面的語句正確的寫法應該是:  float d =  abs( x0 ) + abs(y0)  + abs(z0) - 1.0

 

// Houdini 16.5 Masterclass - OpenCL vs VEX 中的一個案例,作用是改變Volume內點的顏色

 1 kernel void kernelName( 
 2                  int P_length, 
 3                  global float * P ,
 4                  int Cd_length, 
 5                  global float * Cd ,
 6                  int density_stride_x, 
 7                  int density_stride_y, 
 8                  int density_stride_z, 
 9                  int density_stride_offset, 
10                  int density_res_x, 
11                  int density_res_y, 
12                  int density_res_z, 
13                  float16 density_xformtovoxel, 
14                  global float * density 
15 )
16 {
17     int idx = get_global_id(0);
18     if (idx >= Cd_length)
19         return;
20 
21     float3 pos = vload3(idx, P);
22     
23     float4 voxelpos = pos.x * density_xformtovoxel.lo.lo +
24                       pos.y * density_xformtovoxel.lo.hi +
25                       pos.z * density_xformtovoxel.hi.lo +
26                       density_xformtovoxel.hi.hi;
27     int3 voxelidx;
28     voxelidx.x = clamp((int)(floor(voxelpos.x)), 0, density_res_x-1);
29     voxelidx.y = clamp((int)(floor(voxelpos.y)), 0, density_res_y-1);
30     voxelidx.z = clamp((int)(floor(voxelpos.z)), 0, density_res_z-1);
31 
32     float3 c = 1;
33     
34     float d = density[density_stride_offset +
35                       density_stride_x * voxelidx.x +
36                       density_stride_y * voxelidx.y +
37                       density_stride_z * voxelidx.z];
38     if (d > 0.5)
39         c.y = 0;
40         
41     vstore3(c, idx, Cd);
42 }
View Code

 

 需要注意的地方:

1 Volume Transform to Voxel 

勾選這個參數后,在Kernel中會傳遞一個16位矩陣,本例中即  density_xformtovoxel

2 矩陣的乘法

    float4 voxelpos = pos.x * density_xformtovoxel.lo.lo +
                      pos.y * density_xformtovoxel.lo.hi +
                      pos.z * density_xformtovoxel.hi.lo +
                      density_xformtovoxel.hi.hi;

為了得到 voxelpos, 我們用當前點的位置pos 乘以16位矩陣density_formtovoel . pos是一個1x3的矩陣,但在此計算中相當於把他看成 [pos.x,pos.y,pos.z,1] 即1x4的矩陣,乘以4x4的矩陣,這樣的矩陣有好幾種運算方式,這里采用一種分解的簡單方式如下圖。density_xformtovoxel.lo 取得前面一半(8位),density_xformtovoxel.hi取后面一半,density_xformtovoxel.lo.lo取前面的一半的一半(即矩陣的第一行),同理density_xformtovoxel.lo.hi 第二行,density_xformtovoxel.hi.lo第三行,density_xformtovoxel.hi.hi  第四行。

                density_xformtovoxel.lo.lo

density_xformtovoxel = [     density_xformtovoxel.lo.hi      ]

              density_xformtovoxel.hi.lo

                                            density_xformtovoxel.hi.hi 

[ pos.x,  pos.y, pos.z, 1]  x  density_xformtovoxel    根據矩陣的分塊乘法規則可以寫成: 

                      pos.x * density_xformtovoxel.lo.lo +
                      pos.y * density_xformtovoxel.lo.hi + pos.z * density_xformtovoxel.hi.lo + density_xformtovoxel.hi.hi;

3  voxelpos 

voxelpos 得到的是浮點值, clamp 后

1     voxelidx.x = clamp((int)(floor(voxelpos.x)), 0, density_res_x-1);
2     voxelidx.y = clamp((int)(floor(voxelpos.y)), 0, density_res_y-1);
3     voxelidx.z = clamp((int)(floor(voxelpos.z)), 0, density_res_z-1);

我們就可以可到voxel的位置,本例中,Volume分辨率如下

   voxelidx.x  取值范圍 0- 99

   voxelidx.y 取值范圍 0-94

   voxelidx.z 取值范圍 0-109                   

4  density_stride_x, density_stride_y, density_stride_z 

為了理解這三個參數,我在寫入Cd的值前加入三行代碼:

c.x = density_stride_x;
c.y = density_stride_y;
c.z = density_stride_z;

 Volume的分辨率是 [100,95,110]  , 可以看出:

9894 =  (95+2 ) * (100 + 2)

102 = (100 +2)

這是因為當Opencl 拿到我們的volume時, 為了防止訪問volume的數據時超出邊界,會每個方向增加一個voxel作為padding(比如對於X方向,+X軸的最外面和-X軸的最外面各增加一個voxel)。

舉例比如:一個 5X5X3矩陣,當opencl 拿到時,會把它變成7X7X5矩陣,(左邊是原始的volume, 右邊是處理過的)

因此!!!!!!!!!!!!!!

如果我們原始的volume分辨率是 [100,95,110] 

 density_stride_x : x 方向單位增量 是 1 (這個前后不變,加1即可訪問到+X方向的下一個voxel)

density_stride_y   :  y 方向單位增量是  100 + 2 ( 加 102 才能訪問到 +Y方向的下一個voxel)

density_stride_z  :z 方向單位增量   (95+2 ) * (100 + 2) (加9894 才能訪問到 +Z方向的下一個voxel)

             


免責聲明!

本站轉載的文章為個人學習借鑒使用,本站對版權不負任何法律責任。如果侵犯了您的隱私權益,請聯系本站郵箱yoyou2525@163.com刪除。



 
粵ICP備18138465號   © 2018-2025 CODEPRJ.COM