CUDA in Maya plugin

最近一直在想,是否可以在 Maya plugin 裡面使用 CUDA 來作運算。便趁著週末參考 CUDA SDK cppIntegration 的範例,依樣畫葫蘆地在 Maya yTwistNode 的這個範例上作個簡單的小實驗,把 y-twist deform 的計算移至 CUDA kernel function 裡…

環境設定

在撰寫 CUDA 的程式碼之前,我們必須先在原來的 cpp project 中新增一條 build rule…
sshot-vc-cuda

將 CUDA SDK 附的 Cuda.rules 加入列表並啟用…
(為了方便起見,我們可先新增一個環境變數 CUDA_SDK_PATH 使其指向 CUDA SDK 的安裝位置。 而 Cuda.rules 所在的位置則是在 $(CUDA_SDK_PATH)common 裡面)
sshot-vc-cuda

啟用 Cuda.rules 之後, project properties 便會出現一個 CUDA 的分類,一些相關的參數都可以在這裡作設定。 sshot-vc-cuda

程式重構

環境設定完之後,我們新增一個 .cu 的檔案,實作與 cpp code 的溝通介面以及處理 y-twist 的 kernel function

#include
#include 

__global__ void  gpu_twist(float4 pts[], int len,
                           float magnitude, float env)
{
	for(int k=0; k<len; k++)
	{
		float ff = magnitude * pts&#91;k&#93;.y * env;
		if( ff != 0.0f )
		{
			double cct= cosf(ff);
			double cst= sinf(ff);
			double tt = pts&#91;k&#93;.x * cct - pts&#91;k&#93;.z * cst;

			pts&#91;k&#93;.z = pts&#91;k&#93;.x * cst + pts&#91;k&#93;.z * cct;
			pts&#91;k&#93;.x = tt;
		}
	}
}

extern "C"
void	gpu_calc(float ptData&#91;&#93;&#91;4&#93;, unsigned int len,
                        float magnitude, float env)
{
	const unsigned int memSize = len * sizeof(float) * 4;
	float4 *d_data;

	cutilSafeCall( cudaMalloc((void**) &d_data, memSize) );

	// Copy data to gpu memory
	cutilSafeCall( cudaMemcpy(d_data, ptData, memSize,
                                     cudaMemcpyHostToDevice) );

	// Invoke the kernel with one thread
	gpu_twist<<<1,1>>>(d_data, len, magnitude, env);

	// Copy result from gpu memory to host memory
	cutilSafeCall( cudaMemcpy(ptData, d_data, memSize,
                                    cudaMemcpyDeviceToHost) );

	cudaFree(d_data);
}

回到原來的 yTwistNode.cpp 補上溝通介面的宣告,並對 deform function 作一點小幅修改…

extern “C”
void gpu_calc(float ptData[][4], unsigned int len, float magnitude, float env);

MStatus
yTwist::deform( MDataBlock& block,
MItGeometry& iter,
const MMatrix& m,
unsigned int multiIndex)
{
MStatus status = MS::kSuccess;

// determine the angle of the yTwist
//
MDataHandle angleData = block.inputValue(angle,&status);
McheckErr(status, “Error getting angle data handlen”);
double magnitude = angleData.asDouble();

// determine the envelope (this is a global scale factor)
//
MDataHandle envData = block.inputValue(envelope,&status);
McheckErr(status, “Error getting envelope data handlen”);
float env = envData.asFloat();

// { Modified block
MPointArray ptArray;
status = iter.allPositions(ptArray);

// Get raw data
float (*rawData)[4] = new float[ptArray.length()][4];
ptArray.get(rawData);

// Invoke cuda
gpu_calc(rawData, ptArray.length(), magnitude, env);

// Retrieve the data
for(int i=0; i

  • Additional Include Directories
    $(CUDA_INC_PATH); $(CUDA_SDK_PATH)commoninc
  • Additional Library Directories
    $(CUDA_LIB_PATH); $(CUDA_SDK_PATH)commonlib
  • Additional Dependencies
    cudart.lib cutil32D.lib
  • CUDA 的 Additional Include Directories
    $(CUDA_SDK_PATH)commoninc
    Ps. 這裡的 path 記得要 quote 起來,不然會有 nvcc fatal : A single input file is required for a non-link phase when an outputfile is specified 的錯誤訊息!!
  • 最後的步驟當然是 build & execute 囉!! 至於執行效能,還真的如預期中的!? 這是因為 y-twist 其實也沒做多少事,花在 memcpy 的時間搞不好還比計算的時間多。加上我全部的資料都宣告在 global memory , thread 也只用一個而已。所以執行效能較慢也是預料中的事情,不過經過這一次的實驗,總算大致熟悉了從 cpp project 引入 CUDA 的流程,也終於在 CUDA 的世界中心呼喊 Hello World 了!!  😉

    Advertisements

    4 thoughts on “CUDA in Maya plugin

    1. Hello mi:
      如果 pointer 在 device function 是指向 shared memory 或是 global memory 的話,是可以正常運作的。但是在 dereference 的時候,要留意 pointer 指向的位置是在 host 還是 device,如果位置不合法的話就會產生錯誤(像是在 host function 中不小心去 dereference 一個在 device 的位置)。其他的詳情可以參閱 Nvidia CUDA programming guide 的附錄 B.2.5,裡頭有一些相關限制的說明。:)

    Leave a Reply

    Fill in your details below or click an icon to log in:

    WordPress.com Logo

    You are commenting using your WordPress.com account. Log Out / Change )

    Twitter picture

    You are commenting using your Twitter account. Log Out / Change )

    Facebook photo

    You are commenting using your Facebook account. Log Out / Change )

    Google+ photo

    You are commenting using your Google+ account. Log Out / Change )

    Connecting to %s